ck_tile: fix 2:4-sparse SWMMAC correctness on gfx1201/RDNA4 (3 bugs) + fail→pass repro by The-Monk · Pull Request #3759 · ROCm/composable_kernel

The-Monk · 2026-07-24T03:03:01Z

0. TL;DR for the maintainer (schung-amd)

Three correctness bugs in ck_tile's SPARSE MmaOpFamily path produce wrong results on gfx1201/RDNA4 (v_swmmac_*_iu4). Root-caused, fixed (two headers), and reproduced with a committed standalone test that fails on current CK and passes with the fix. We found these while building a production-grade 2:4-sparse fp8/int4 GEMM for LLM inference on RDNA4 — §3-4 give that context, which doubles as validation that the fixed path works end-to-end at scale.

1. The three bugs (files: `sparse_mma_pipeline.hpp`, `sparse_transforms.hpp`)

Bug 1 — compress_a_impl writes phantom values on <2-nonzero 2:4 groups.
The fallback ADataType nonzero_elems[2] = {a_vec[i*4+2], a_vec[i*4+3]} seeds the compressed pair from fixed input positions rather than true zero. A 2:4 group with 0 or 1 real nonzeros then reconstructs with a phantom value from a wrong lane (observed as register aliasing / UB). Fix: default nonzero_elems to true {0,0} and fill only real survivors — 0/1/2-nonzero groups all reconstruct exactly.

Bug 2 — packed sub-byte (pk_int4_t) group-of-4 scan treats a byte as one element.
The compaction's "nonzero per group of 4" scan operates on bytes, but a pk_int4_t byte holds two 4-bit values; a byte is "nonzero" if either nibble is, so the group-of-4 assumption over-counts and writes out of bounds into nonzero_elems[2]/[3]. Fix: packed-sub-byte-aware nibble counting (guarded with a static_assert that the packed path currently implements only pk_int4_t).

Bug 3 — compress_a_impl emits the two per-group idx metadata fields in an order that mismatches what the hardware reads, requiring a SWAP + XOR-1 to correct.
Within CK's own pk_int4_t packing, the metadata CK writes has the two idx fields swapped relative to which compressed survivor they govern, plus a XOR-1, versus what v_swmmac_i32_iu4 consumes on gfx1201 — so the sparse result comes out wrong (pos(HIGH-nibble survivor, found first) = idx1_field XOR 1, pos(LOW-nibble survivor, found second) = idx0_field XOR 1). Dense iu4 was confirmed correct first (11 seeds) to isolate this to the sparse metadata path. Fix: emit the idx fields with the swap+XOR so the hardware reconstructs the correct positions (Idx0 < Idx1, the documented 2:4 contract, is naturally satisfied).
Precision note (so this isn't over-stated): we verified this is specific to CK's compress/packing convention, not a universal "the hardware swaps." An independently-derived encoder in our own driver produces correct metadata without the swap — and applying the swap to it breaks a working encoding. So the fix corrects CK's own ordering to match the hardware; it is not a claim about v_swmmac_iu4 behavior for all encoders.

2. The test (committed; fails on current CK, passes with the fix)

test/ck_tile/gfx1201_sparse_swmmac/sparse_swmmac_correctness_repro.cpp drives CK's own machinery for the SPARSE MmaOpFamily — it builds Pipeline::AWarpDstrEncoding internally (no manual compression math), generates an adversarial 2:4 A tile in-source (each 4-group cycles through all six two-survivor position pairs, all four single-survivor positions, and the zero-survivor group — survivors at every position, values keyed to (row, group, position) so misplacement shows numerically rather than cancelling), runs the SWMMAC GEMM (K=32/64/128, three FragsK tile sizes), and checks against a CPU int64-accumulate reference. The legacy slots-0,2 pattern is retained as an opt-in control (-DUSE_CANONICAL_PATTERN); it cannot detect Bug 1 and passes on both trees.

Verification (2026-08-15, head ad79359d0 vs merge-base 8fc1ac24e9, identical test source on both trees, 1x R9700 gfx1201, ROCm 7.14):

Unfixed (merge-base): FAIL — max_abs_err = 112 / 127 / 272 at K=32/64/128
Fixed (this PR): PASS — max_abs_err = 0 on all three shapes

Reproduction (self-contained — no external files):

git clone -b gfx1201-sparse-swmmac-fixes https://github.com/The-Monk/composable_kernel ck-3759 && cd ck-3759
hipcc -std=c++17 -O2 --offload-arch=gfx1201 -I include \
  test/ck_tile/gfx1201_sparse_swmmac/sparse_swmmac_correctness_repro.cpp -o repro
HIP_VISIBLE_DEVICES=0 ./repro     # -> ALL PASS (max_abs_err=0)
# unfixed A/B: git worktree add ../ck-base $(git merge-base HEAD origin/develop)
#   then build the SAME test file with -I ../ck-base/include -> FAIL (112/127/272)

Historical note: earlier revisions of this description cited max_abs_err=352 from a real Quark-quantized weight tile (REAL_A_16x128, -DUSE_REAL_TILE); that input is superseded by the in-source generator (same root cause, now reproducible from the PR alone) — the 352 figures apply only to that tile.
(We compiled with the ROCm-devel clang toolchain directly — clang++ -x hip --offload-arch=gfx1201; the distro hipcc on this box is a stale 5.7-era wrapper unrelated to CK's target toolchain and will not build ck_tile headers. Any current ROCm/HIP clang works the same way.)

ASAN attempted, not obtained (honest note, not a blocker): we tried to get a device-sanitizer stack trace on the Bug-1/Bug-2 OOB write for extra evidence (-fsanitize=address -fgpu-sanitize). Device ASAN needs an xnack+ target-ID variant; RDNA4 (gfx1201) does not accept an xnack+/xnack- feature suffix at all (clang++: error: invalid target ID 'gfx1201:xnack+') — HIP device-side ASAN is a gfx9/CDNA-class feature (gfx90a/gfx94x with xnack+), not available on RDNA4 consumer/workstation targets in this ROCm toolchain. We did not chase this further since it's not required to reproduce the bug — the max_abs_err numbers above are a complete, deterministic fail→pass.

We'll port this into CK's gtest format for the PR; the standalone repro is included so it's runnable without the full CK test build.

3. Context — why these bugs mattered (the "whole thing")

We hit these building a library-grade 2:4-sparse GEMM for RDNA4 LLM inference (llama.cpp/ggml fork, gfx1201, 2× R9700). The journey, honestly:

ISA ceiling is real: isolated v_swmmac microbench hits 765 TOP/s fp8-2:4 (2× dense) and 1531 TOP/s int4-2:4 (4×) at ILP≥4 — 88-100% of the R9700 spec.
A hand-written MMQ-grade 2:4 kernel (cooperative tiles, LDS staging, ILP≥4) reaches ~93% of our native-fp8 WMMA MMQ (which we built — none exists upstream) on the whole model, and is at or ahead of the library on the three dominant GEMM shapes (+1.8% / +3.9% / +32% on gate-up / q-o / down-proj).
The honest finding worth AMD seeing: the ISA 2×/4× sparsity ceiling does not translate to a compute-bound, well-tiled model GEMM — measured across three quality tiers, the sparsity factor is 2× (isolated ISA) → 1.25× (crude kernel) → ~1.01× (library-grade). Once the kernel is well-tiled, the GEMM is bound by the tiled system (LDS staging, occupancy), not the raw SWMMAC issue rate that sparsity accelerates. This is a real, reproducible microbench-vs-model boundary for RDNA4 structured sparsity.
The mechanism, made precise (int4-2:4): the "4×" decomposes as 2×(int4 vs int8) × 2×(2:4 sparsity). The first 2× is real and already captured by dense wmma_i32_16x16x32_iu4 — gfx1201's dense int4 engine is already K=32-wide. The second 2× washes: swmmac_i32_16x16x32_iu4 (749,761 GOP/s) ≈ wmma_i32_16x16x32_iu4 (765,061) = 0.98×, instruction-for-instruction, because dense int4 already occupies the K=32 slot. The only sparse tensor edge is the K=64 form (1.77×) — the same "2×-K" mechanism fp8-2:4 already showed washing at library grade. So on high-arithmetic-intensity LLM GEMM shapes (compute-bound), structured 2:4 is a ~1.2× weight-bandwidth win (memory-bound regime only), not the tensor 2×/4×. (Correctness first regardless — the perf question is separate, and now answered.)
The remaining gap to the library was pinned by direct profiling to two specific, non-magical things — a 2×-slow activation-quantize helper (LDS+sync vs shuffle) and one occupancy-starved small-N shape — both fixed [§4].

4. Closing the gap — result

The 8% was pinned (by direct rocprofv3 head-to-head) to two things, NOT the big GEMMs (we're at/ahead of our own fp8 MMQ there: +1.8% / +3.9% / +32% on gate-up / q-o / down-proj). Fixing them:

build	dense-MMQ (vs our fp8 MMQ)	2:4-MMQ (vs our fp8 MMQ)
capstone baseline	3786 (0.924×)	3823 (0.933×)
+ shuffle-based quantize (kept)	3900 (0.952–0.978×)	3944 (0.962–0.989×)

Quantize helper rewrite (kept): replaced a 32-thread / LDS+__syncthreads() reduction with the upstream mmq.cuh quantize shape — 128 threads, float4 loads, warp-shuffle max-reduction, zero LDS/sync (halved it, 36µs→~18µs). Closed ~40% of the remaining gap; both kernels now high-90s% of our native-fp8 kernel (0.989× same-session).
The N=1024 k/v small-shape (honest negative): both candidate levers — LDS bank-conflict padding and coarser K-per-sync — were implemented, correctness-verified, and measured worse (padding −2.5%; coarser-K a monotonic regression as it forced dropping double-buffering into an occupancy/LDS cliff). Both reverted. The outlier remains open, wants a different lever (shape-adaptive smaller tile without touching sync granularity).

Net: the hand-written 2:4 kernel now reaches ~96–99% of our native-fp8 WMMA MMQ (which we built — none exists upstream), is ahead on the FLOP-dominant GEMMs, and carries the measured ~1% sparsity edge over a comparable dense kernel — a library-competitive RDNA4 2:4 GEMM on top of the correctness fixes.

5. A related RDNA4 iu4 quirk (same class, different instruction — FYI, not part of the fix)

The Bug-3 element-ordering quirk on swmmac_iu4 is not isolated: an independent kernel-authoring effort on this box hit the same class on the dense v_dot8_i32_iu4/sudot8 path — "dots mismatched element pairs." RDNA4's iu4 instruction family appears to carry undocumented element-ordering conventions in both the sparse (SWMMAC metadata) and dense (dot8 operand pairing) paths. Documenting these in the ISA/CK would save the next implementer the multi-day reverse-engineering we did. Happy to write up the dense one too if useful.

6. Method (for reproducibility)

All numbers: gfx1201 (R9700), ROCm 7.14, GPU-isolated, warm, medians ≥3. Correctness gated by execution (CPU-reference compare), never inspection. The full capability-optimizer method (measure → grade vs the published peak → drive the lever) is what surfaced both the bugs and the microbench-vs-model boundary.

[CK][CK TILE] Clean up tile_engine grouped_conv harness (#7761) ## Motivation Tile_engine grouped_conv contains ML heuristic validation scripts which cause confusion to new developers. So, this PR is intended to relocate the scripts into dispatcher/heuristic directory to maintain separation of concern. ## Technical Details The grouped_conv tile_engine directory is a benchmarking harness for grouped convolution kernels; ML-heuristic content does not belong there. - Move compare_ml_vs_oracle.py and validate_ml_vs_oracle.py from tile_engine/ops/grouped_conv/ to dispatcher/heuristics/validation/grouped_conv/, and rebase their sys.path / oracle CSV / model dir lookups for the new location (CSV path is now an --oracle-csv flag instead of a hard-coded sibling). - Move GROUPED_CONV_HEURISTIC_REPORT.md (system-level ML report) into dispatcher/heuristics/ where the rest of the heuristic docs live. - Rewrite tile_engine/ops/grouped_conv/README.md as a pure benchmarking / dispatcher-sweep doc (kernel enumeration, JIT pipeline, CSV schema, problem registry), in the style of tile_engine/ops/fmha/README.md. All ML training / model-efficiency content is removed and replaced with a pointer to dispatcher/heuristics/. ## Test Plan Validation scripts are re-wired and tested locally ## Test Result Tests passed on local machine. ## Submission Checklist - [x ] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

ck_tile: add FillUniformScaleDistribution and fix MX GEMM scale init (#7724) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit ## Summary ### Problem MX GEMM pipeline tests were passing vacuously: scale bytes were drawn from a fixed range (40–60) which, for e8m0, maps to scales ≈ 10⁻²⁷ — far below FP16 min denorm. Both GPU and CPU produced all-zero outputs, so numerical checks passed without exercising the GEMM. ### Changes **`include/ck_tile/host/fill.hpp`** — new `FillUniformScaleDistribution<ScaleType>` functor - Accepts human-readable float bounds and maps them to the raw byte range of any ExMy scale type (e8m0, e4m3, e5m3) by re-centering the IEEE 754 exponent into the type's bias space - Sampling is uniform over raw bytes → uniform over representable values - Fixes left-shift UB: uses multiplication instead of `<< mant_bits` to avoid shifting negative signed integers (C++17 UB) - Adds `assert(min_r <= max_r)` to catch inverted-range UB when both bounds exceed the type's representable range - Provides default member values (0.125f, 2.0f) and `std::optional` seed consistent with sibling fillers - `/** */` Doxygen style with `@note` on snapping asymmetry **`test/ck_tile/gemm_mx/test_mx_gemm_pipeline_util.hpp`** — fix scale initialization - Replace manual byte-range distribution with `FillUniformScaleDistribution<>{0.125f, 2.0f}` - Use distinct seeds for scale_a (11941) and scale_b (11943) to avoid correlated scale tensors that were causing 60 test failures for fp4+e5m3/e4m3 combinations **`test/ck_tile/utility/test_fill.cpp`** — new unit tests for `FillUniformScaleDistribution` - 16 typed tests across e8m0, e4m3, e5m3: validity, range, reproducibility, coverage, snapping, stress, nullopt seed, and range overload - Test helper `expected_raw_range` mirrors implementation clamping exactly

[CK] add credentials to docker manifest inspect call ## Motivation This should fix an issue that we recently encountered in CI when we exceeded the limit of accessing docker without authentication: [2026-05-29T16:08:42.447Z] + docker manifest inspect --insecure rocm/composable_kernel:ck_ub24.04_rocm7.13 [2026-05-29T16:08:42.833Z] toomanyrequests: You have reached your unauthenticated pull rate limit. https://www.docker.com/increase-rate-limit ## Technical Details  ## Test Plan  ## Test Result  ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

[CK_TILE] Fix Stream-K k_size calculation ## Motivation In a recent benchmarking task for CK Tile Stream-K algorithm, we identified that certain instances segfault. This change works to fix the bug and adds necessary regression tests. ## Technical Details The StreamK kernel constructs tensor views using a `k_size` parameter that determines how much of the K dimension to process in each iteration. Previously, this was calculated as: ```cpp index_t k_size = num_loop_sk * TilePartitioner::KPerBlock; ``` This calculation assumes all macro tiles along K are exactly `KPerBlock` in size. However, when `K % KPerBlock != 0`, the final macro tile along K has a remainder size of `K % KPerBlock`, not a full `KPerBlock` (see the figure below): <img width="961" height="488" alt="image" src="https://github.com/user-attachments/assets/3e1cceed-5dcd-4980-8b02-cee24eecf262" /> With the old code, a workgroup working with the `MPerBlock x (K % KPerBlock)` tile in A and B risk accessing illegal memory. Hence, this change ensures that when `K % KPerBlock != 0`, workgroups processing iterations that include the final macro-tile along K calculate the correct `k_size` based on the remainder rather than assuming a full `KPerBlock`. ## Test Plan I added the following tests: 1. Unit tests added for the Stream-K Tile Partitioner: - `StreamKTilePartitionerBaseGetKSize/NoRemainderTiles` - validates full tiles - `StreamKTilePartitionerBaseGetKSize/RemainderTiles` - validates remainder handling 2. Regression tests that test a case where `K % KPerBlock != 0` ## Test Result Tests passed locally on gfx90a, gfx942, and gfx950. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

[CK_TILE] Use gfx11 float buffer atomics in FMHA Bwd ## Motivation FlashAttention CK backward on gfx11 can hit out-of-bounds/tail writes in the dQ accumulator atomic-add path when sequence rows are padded at the tile level but not marked invalid in the DQDKDV main tensor view. With the generic global atomic fallback, an incorrectly-valid tail element can issue an actual pointer-based `atomicAdd`. With the buffer atomic path, the write is issued through a buffer resource with bounds information and follows the same backend already used by gfx9/gfx12. This fixes the gfx11 FMHA BWD failure without changing the gfx11 default for unrelated CK Tile kernels. ## Technical Details This PR enables the existing CK Tile AMD buffer float atomic-add path only for generated FMHA BWD gfx11 translation units. gfx11 normally uses the generic global atomic fallback for floating-point `buffer_view::atomic_add`. That fallback performs the atomic through a raw computed pointer and depends on the software validity predicate to avoid invalid elements. In FMHA BWD dQ accumulation, padded tail rows can reach this path, so using the buffer atomic backend is safer: it uses a buffer resource with base pointer, bounds information, and an element offset, matching the backend already used by gfx9/gfx12. Enabling `CK_TILE_USE_AMD_BUFFER_ATOMIC_ADD_FLOAT` globally for gfx11 is too broad and can break unrelated gfx11 CK builds such as GEMM. Instead, `config.hpp` now preserves an explicitly pre-defined `CK_TILE_USE_AMD_BUFFER_ATOMIC_ADD_FLOAT`, while keeping the existing default disabled for gfx11. ## Test Plan Validated the change with the FlashAttention CK full test suite with backward pass enabled on gfx11. pytest -q -s tests/test_flash_attn_ck.py ## Test Result FlashAttention CK gfx11 test result: 260680 passed, 152076 skipped ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. Co-authored-by: Po Yen Chen <PoYen.Chen@amd.com>

[CK] apply the compiler warning suppression flags in cmake files (#7863) ## Motivation Apply the blanket suppression flags for latest clang warnings in staging compiler such as: lifetime-safety-lifetimebound-violation lifetime-safety-intra-tu-suggestions lifetime-safety-cross-tu-suggestions unknown-warning-option ## Technical Details  ## Test Plan  ## Test Result  ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

[CK_Tile] Add scale16 Support for F4 WMMA in CK_Tile ## Motivation This PR adds CK Tile support for the scale16 F4 WMMA path on gfx1250 and improves warp GEMM unit test coverage/structure for gfx1250-specific cases. ## Technical Details - Scale16 support in warp GEMM dispatch and WMMA trait plumbing: added IsScale16 plumbing to warp GEMM dispatcher path - Warp GEMM test restructuring for gfx1250: added Warp GEMM gfx1250 coverage to verify all F4 WMMA paths ## Test Plan Run ./test_ck_tile_wg_32x16x128_fp4. ## Test Result ``` ./test_ck_tile_wg_32x16x128_fp4 [----------] Global test environment tear-down [==========] 3 tests from 1 test suite ran. (1751 ms total) [ PASSED ] 3 tests. ``` ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

[CK_TILE] Fix conditional rescale numerical instability in FMHA forward (#6498) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit [CK_TILE] Fix conditional rescale numerical instability in FMHA forward ## Motivation Fix numerical instability in the conditional O-accumulator rescaling optimization for CK-Tile FMHA forward (FlashAttention-4, Algorithm 6, Eq. 6). The conditional rescale optimization skips the expensive O-accumulator rescale when the running row-max shift is within a threshold (tau = log2(256) = 8.0). The original implementation had a bug: attention weights P were computed in the `m_new` reference frame before the skip/rescale decision. In the skip branch, `m` was reverted to `m_old`, but P remained in the `m_new` frame, causing incorrect softmax normalization. This fix introduces a `p_row_correction` factor: in the skip branch, P is multiplied by `exp2(m_new - m_old)` to bring it back to the `m_old` reference frame. - **Correctness:** Fixes broken inference on long sequences where running-max drift causes exp2 overflow (observed as degraded image quality on MI350X Flux2 generation) - **Performance:** Neutral to +4% depending on workload shape ## Technical Details 6 pipeline header files (same pattern in each): - `block_fmha_pipeline_qr_ks_vs.hpp` - `block_fmha_pipeline_qr_ks_vs_async.hpp` - `block_fmha_pipeline_qr_ks_vs_async_trload.hpp` - `block_fmha_pipeline_qr_ks_vs_fp8.hpp` - `block_fmha_pipeline_qr_ks_vs_whole_k_prefetch.hpp` - `block_fmha_pipeline_qs_ks_vs.hpp` In each file: - Lower threshold from 10.0 to 8.0 (tau = log2(256)) - Add `p_row_correction` distributed tensor initialized to 1.0 - Rescale branch: standard rescale of O_acc and l; correction = 1.0 - Skip branch: compute correction = exp2(-acc_scale_log2), update l, revert m, store correction - New `p_spans` sweep applies per-row correction to `p_compute` before P*V GEMM - Move P-to-PDataType cast to after correction sweep ## Dependencies None — this PR is standalone. ## Test Plan - GPU validation on MI300X (gfx942, ROCm 6.4.1): - Command: `./build/bin/tile_example_fmha_fwd -b=2 -h=8 -s=4096 -d=128 -prec=bf16 -v=1 -warmup=1 -repeat=3` - GPU validation on MI350X (gfx950, ROCm 7.0): - Command: `./build/bin/tile_example_fmha_fwd -b=2 -h=8 -s=4096 -d=128 -prec=bf16 -v=1 -warmup=1 -repeat=3` - Command: `./build/bin/tile_example_fmha_fwd -b=2 -h=8 -s=4096 -d=128 -prec=fp16 -v=1 -warmup=1 -repeat=3` ## Test Result Accuracy vs FP32 reference (MI350X, gfx950): | Shape | max_diff | mean_diff | |-------|----------|-----------| | B=1 H=24 M=4096 K=128 bf16 | 9.1e-4 | 4.6e-5 | | B=4 H=32 M=4096 K=128 bf16 | 9.9e-4 | 4.6e-5 | | B=1 H=24 M=4096 K=128 fp16 | 1.2e-4 | 9.0e-6 | Performance (MI350X, gfx950, ROCm 7.0): | Shape | FA4 (TFlops) | Always-rescale (TFlops) | Delta | |-------|-------------|------------------------|-------| | B=1 H=24 M=4096 K=128 bf16 | 425.9 | 428.5 | neutral | | B=2 H=8 M=2048 K=256 bf16 | 513.9 | 509.0 | +1.0% | | B=1 H=64 M=2048 K=64 bf16 | 481.7 | 464.3 | +3.7% | Benchmark results (MI300X, gfx942, ROCm 6.4.1): No regression on MI300X. This correctness fix is performance-neutral. | Config | TFlops / GB/s | Time (ms) | |--------|-------------|-----------| | MHA bf16 b=2 h=8 s=4096 d=128 | 342.49 TFlops | 0.401 | | MHA fp16 b=2 h=8 s=4096 d=128 | 391.70 TFlops | 0.351 | | Causal MHA bf16 b=2 h=8 s=4096 d=128 | 227.07 TFlops | 0.303 | | GQA 4:1 bf16 b=2 h=32 hk=8 s=2048 d=128 | 324.69 TFlops | 0.423 | | GQA 8:1 bf16 b=2 h=64 hk=8 s=2048 d=128 | 348.09 TFlops | 0.790 | | LLaMA-70B prefill b=1 h=64 hk=8 s=4096 d=128 bf16 | 376.71 TFlops | 1.459 | | Long-seq bf16 b=1 h=16 s=16384 d=128 | 383.42 TFlops | 5.735 | | Decode b=64 h=32 hk=8 s_k=4096 d=128 bf16 | 691.64 GB/s | 1.554 | All validation tests pass (`valid:y`) on both MI300X and MI350X. Additional validation: - Uniform scores: softmax output matches FP32 reference (max_diff < 1e-3) - Large seqlen (4096+): no overflow or NaN in O-accumulator - Spike pattern: correct handling of sudden row-max jumps - Multiple spikes: correction applied correctly across multiple skip/rescale transitions - Deterministic: identical outputs across repeated runs - No performance regression on standard workloads

[CK] Extract Jenkinsfile helpers into vars/ck.groovy shared library (#7743) ## Motivation The CK Jenkinsfile is a 2,215-line monolith mixing helper function definitions with pipeline stage declarations. This makes it difficult to review, modify, or extend CI stages without wading through unrelated infrastructure code. ## Technical Details Extract all helper functions from the Jenkinsfile into vars/ck.groovy, loaded at runtime via ck = load "vars/ck.groovy" in the first stage. The Jenkinsfile is reduced from 2,215 lines to 810 lines containing only the pipeline structure. - 36 helper functions moved to ck.groovy with no logic changes - 10 new stage-wrapper functions (runBuildCKAndTests, runTileEngineGemmTests, runClangFormat, etc.) extract inline environment{}/steps{} business logic from stages, eliminating the MethodTooLargeException caused by CPS-transformed shell strings exceeding the JVM 64KB bytecode limit - All ck. method calls in steps{} blocks wrapped in script{} as required by Jenkins Declarative Pipeline - rocmnode() remains in the Jenkinsfile (needed for agent{} labels before ck is loaded) - CRON_SETTINGS / POLL_SPEC remain in the Jenkinsfile (triggers{} evaluates at parse time before any workspace is available) - No stage names changed ## Test Plan - Jenkinsfile validated against the Jenkins Pipeline Linter (/pipeline-model-converter/validate) - All 35 shared helper functions diffed line-by-line against develop to verify no regressions - Merge from develop incorporated and verified (gfx1250 stage, ROCm 7.13 default, cmake_build updates) ## Test Result - Linter: passes - Function diff vs develop: all 35 functions match exactly - Awaiting Jenkins run to confirm end-to-end stage execution ## Submission Checklist - [ x ] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

[CK Tile] Add conv Wavelet GEMM pipeline and bwd_weight instances (#7937) ## Motivation CK Tile had no pipeline competitive with old CK's wavelet on the RetinaNet K=36 C=256 3x3 conv bwd_weight class. This adds a wave-specialized "wavelet" GEMM pipeline so CK Tile has a competitive kernel for spatial small-K shapes. ## Technical Details - New wavelet GEMM pipeline (`gemm_pipeline_ag_bg_cr_wavelet.hpp`): workgroup split into math waves (LDS read + MFMA) and load waves (DRAM read + LDS write). - VGPR role-split: `operator()` has two top-level mutually-exclusive `is_math` branches so the allocator overlays both roles onto the same physical VGPRs, cutting arch VGPR ~33-40% and raising occupancy. Correctness depends on identical `block_sync_lds` counts on both arms plus a matching load-wave barrier stub in the epilogue (`cshuffle_epilogue.hpp`). - Kernel dispatch (`grouped_convolution_backward_weight_kernel.hpp`): `kIsWavelet` path, `LaunchBlockSize`, load-wave barrier stub. Uplift: wavelet is the fastest CK Tile pipeline on the RetinaNet K=36 C=256 3x3 family, beating the best non-wavelet CK Tile kernel by 10-27% (googlenet K=320 by 16-23%); the role-split roughly halves the parity gap vs old CK on the 13x13 fp16 shape. ## Test Plan - `ckProfiler grouped_conv_bwd_weight`, NHWGC layout, fp16/bf16, `split_k=all`, CPU verify on RetinaNet K=36 shapes (7x7, 13x13) and a broad 2D sweep. - Correctness: `-v=1` across `split_k` in {-1,1,2,4,8,16,32,64} (barrier-parity / deadlock check). - `test_grouped_convnd_bwd_weight` over the tests `.conf` wavelet instances. ## Test Result - All wavelet instances CPU-verify correct across the split-K sweep; no hangs (dual-arm barrier sequence matches). - Wavelet wins the RetinaNet K=36 C=256 3x3 family (10-27% over best non-wavelet CK Tile) and googlenet K=320 (16-23%); at parity-or-better vs old CK on the majority of spatial shapes. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

[CK] Allow skipping split-K C-buffer zero-init in xdl_cshuffle blockscale GEMM (#7935) Add a `skip_zero_init` flag (default false) to the Problem/Argument of the xdl_cshuffle block-scale GEMM device ops (multiple_d ab_scale and blockscale b-preshuffle). When the flag is set, the device invoker skips the internal hipMemsetAsync that zeroes p_c_grid before the KBatch > 1 split-K atomic-accumulation path. The flag is declared on the gridwise Problem struct (inherited by Argument), so it is visible on both the rotating-cache (arg_) and the normal (arg) launch paths in each device op. Why: callers that already pre-zero the output buffer otherwise pay for a redundant device-wide memset before split-K atomic accumulation. Gating the memset behind an opt-in flag lets such callers avoid the duplicate work. Because the flag defaults to false, every existing call site is unaffected and the observable behavior is unchanged. ## Motivation  ## Technical Details  ## Test Plan  ## Test Result  ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. Co-authored-by: Cursor <cursoragent@cursor.com>

[CK_Tile][MI450] Add bf16 output wmma instruction (16x16x32) (#7830) Wire __builtin_amdgcn_wmma_bf16_16x16x32_bf16 into CK Tile for gfx1250, enabling bf16-input bf16-output WMMA at the warp GEMM level. - Add WmmaTraits specialization for <gfx125_t, bf16, bf16, bf16, 16,16,32> - Add WarpGemmAttributeWmmaImpl typedef and WarpGemmWmma alias - Add Dispatcher entry for bf16->bf16 16x16x32 - Add warp_gemm test with reference GEMM validation ## Motivation  ## Technical Details  ## Test Plan  ## Test Result  ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

Replace nested static_for lambdas with compile-time search helper (#6696) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit ## Summary - Add `sequence_find_value` and `find_in_tuple_of_sequences` compile-time search helpers with O(1) template depth - Replace nested `static_for` lambdas in `TensorDescriptor::GetTransformAndItsUpperDimension` and `InitializeElementSize` - Apply same optimizations to `TensorAdaptor` Supersedes #4287. Conflict-resolved rebase of ROCm#3600 onto current develop. ## Motivation The `TensorDescriptor` and `TensorAdaptor` classes had excessive template instantiation from: 1. Nested `static_for` loops with lambdas creating unique closure types at every call site 2. `generate_tuple` with lambdas causing per-type instantiation overhead The new helpers use constexpr array lookup and pack expansion instead of recursive template patterns, achieving O(1) template depth. ## Results (`example_grouped_conv_fwd_xdl_fp16`, n=10, interleaved, `-j1`, `-ftime-trace`) | TU | Baseline (mean) | New (mean) | Delta | Wilcoxon p | Mann-Whitney p | |----|-----------------|------------|-------|-----------|---------------| | `grouped_conv_fwd_xdl_fp16` (host) | 14,886 ms | 13,353 ms | **-10.3%** | **0.002** | **0.0002** | | `grouped_conv_fwd_xdl_fp16` (device) | 27,762 ms | 25,629 ms | **-7.7%** | **0.002** | **0.0002** | | **Total (all TUs)** | **57,732 ms** | **54,030 ms** | **-6.4%** | | | Unrelated TUs (`device_memory`, `host_tensor`, `convolution_parameter`) show no significant difference (p > 0.3), serving as negative controls. ### Methodology - 10 interleaved runs (baseline₁, new₁, baseline₂, new₂, ...) on the same node to eliminate ordering/warmup bias - Wilcoxon signed-rank test (paired, non-parametric) and Mann-Whitney U test (unpaired) - Built with patched clang (LLVM 22) on ctr2-alola-compile-11, `-j1` for accurate per-TU timing - Raw data available in Slurm job 275230 results ## Test plan - [x] 11 unit tests added (5 for `sequence_find_value`, 6 for `find_in_tuple_of_sequences`) - [x] Compile-time benchmark with statistical significance (p < 0.01) - [ ] Full CI Tracking issue: #4229

[CK] Upgrade to new gfx1250 compiler and fix build issues (#7960) ## Motivation The docker image we've been using to build for gfx1250 is a few months old, so we need to upgrade. Some of the changes in the latest compiler version require changes in the code. TDM is temporarily disabled due to changes in the lds load/store intrinsics. ## Technical Details  ## Test Plan  ## Test Result  ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

[CK] Fix gfx950 AITER Sync Regressions MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit ## Summary Fixes three gfx950 regressions in the AITER downstream CI that surfaced after the internal/gfx1250 re-sync (ROCm/rocm-libraries#6978): > **Companion aiter PR:** ROCm/aiter#3392 — host-side adaptations (`Kernel::BlockSize()` `constexpr` drops, blockscale `KBatch=1` clamp) plus the CK submodule bump used to validate these fixes together. - **FlyDSL MoE AOT cache miss** — the AITER MoE tests run with `check_aot_cache=True` and fail on any FlyDSL JIT cache miss, but the CI never pre-compiles the FlyDSL MoE kernels, so gfx950 always misses. Pre-compile them at the start of the AITER test stage. - **`buffer.load.lds.v4i32` link error** — ROCm/rocm-libraries#6978 reintroduced a clang-version guard mapping `llvm.amdgcn.raw.buffer.load.lds` to a `.v4i32`-suffixed name. That name exists in no LLVM (the rsrc operand is a fixed, non-overloaded `<4 x i32>`, so the intrinsic is never type-mangled), so gfx950 4-DWORD direct-to-LDS (e.g. fp4 MoE bpreshuffle) fails to link with `lld: undefined symbol: llvm.amdgcn.raw.buffer.load.lds.v4i32`. Use the canonical plain name unconditionally. - **mixed-precision flatmm warp-GEMM call** — ROCm/rocm-libraries#6978 generalized the scaled `WarpGemmImpl::operator()` from a fixed `<index_t opselA, index_t opselB>` signature to a variadic `<typename... Params>` one and updated the `mx_flatmm` pipeline to pass the op-selectors as `OpSelA<>`/`OpSelB<>` types, but missed the mixed-precision flatmm pipeline (`F8xMXF4`/`F16xMXF4`), which still passed raw integer op-selectors. These no longer bind to `typename... Params` (`error: no matching member function for call to 'operator()'`), breaking compilation of the fp8/bf16 × fp4 cktile MoE gemm1 instances on gfx950 (aiter `test_moe_2stage`). Wrap the op-selectors in `OpSelA<>`/`OpSelB<>`. ## Changes - `Jenkinsfile`: pre-compile the FlyDSL MoE AOT cache (`python3 aiter/aot/flydsl/moe.py`) before the AITER tests. - `include/ck/utility/amd_buffer_addressing_builtins.hpp` and `include/ck_tile/core/arch/amd_buffer_addressing_builtins.hpp`: drop the `__clang_major__` guard and always use `__asm("llvm.amdgcn.raw.buffer.load.lds")`. The plain name is the canonical one for all sizes including the gfx950 16-byte form, as the upstream LLVM gfx950 tests confirm. - `include/ck_tile/ops/flatmm/pipeline/mixed_prec_flatmm_pipeline_agmem_bgmem_creg_v1.hpp`: wrap the warp-GEMM op-selectors in `OpSelA<>`/`OpSelB<>` at the five call sites, matching the `mx_flatmm` pipeline. ## Test plan Validated via CI.

[CK_TILE][FMHA] Optimize long-context decoding on gfx11/12 (#7500) ## Motivation Relevant issue: ROCM-22065 FMHA has less-than-optimal performance of long-context decoding (i.e. when seqlen_q = 1) on gfx11/12. This PR optimizes the splitkv pipeline and configs for such scenarios. ## Technical Details Optimizations applied in this PR: 1. use tiles with smaller M0 (16 vs 64), these tiles are used when seqlen_q <= 16 2. adapt qr_nwarp_sshuffle pipeline for gfx11, it allows to use more warps even for M0 = 16 (the qr pipeline parallelizes work between warps in M dim so with M0 = 16 it allows to use only 1 warp) 3. enable kMergeNumHeadGroupsSeqLenQ (an optimization that merges one group of heads in GQA) for all hdim values, not only 128 4. increase the number of splits (multiply by the number of head groups) if (3) is used 5. increase the number of splits for RDNAs (`multiProcessorCount` is the number of WGPs on RDNAs, not CUs, so it should be doubled to have meaning similar to CDNAs) Performance on gfx1151: | Case | develop (GB/s) | This PR (GB/s) | |:-------|-------:|-------:| | [fp16\|group\|bshd] b:1, h:32/32, s:1/45056, d:64/64 | 127.58 | 183.11 | | [fp16\|group\|bhsd] b:1, h:32/32, s:1/45056, d:64/64 | 153.64 | 215.02 | | [fp16\|group\|bshd] b:1, h:16/8, s:1/77184, d:128/128 | 120.51 | 225.76 | | [fp16\|group\|bhsd] b:1, h:16/8, s:1/77184, d:128/128 | 130.62 | 223.84 | | [fp16\|group\|bshd] b:1, h:32/32, s:1/9600, d:128/128 | 82.65 | 138.44 | | [fp16\|group\|bhsd] b:1, h:32/32, s:1/9600, d:128/128 | 105.75 | 220.45 | | [fp16\|group\|bshd] b:1, h:8/1, s:1/401024, d:256/256 | 16.27 | 187.89 | | [fp16\|group\|bhsd] b:1, h:8/1, s:1/401024, d:256/256 | 16.28 | 188.19 | ## Test Plan An additional test case is added to the exiting test. It uses seqlen_q = 1, GQA, no mask to trigger the changes ``` ninja test_ck_tile_fmha_fwd_fp16 && bin/test_ck_tile_fmha_fwd_fp16 --gtest_filter="*SplitKV* ninja test_ck_tile_fmha_fwd_bf16 && bin/test_ck_tile_fmha_fwd_bf16 --gtest_filter="*SplitKV* ``` Manual testing can be done with these commands: ``` bin/tile_example_fmha_fwd -prec=fp16 -mode=1 -page_block_size=128 -b=1 -h=32 -h_k=32 -d=64 -s=1 -s_k=$((352 * 128)) -lse=1 -mask=0 -num_splits=0 -kname=1 -v=1 bin/tile_example_fmha_fwd -prec=fp16 -mode=1 -page_block_size=128 -b=1 -h=16 -h_k=8 -d=128 -s=1 -s_k=$((603 * 128)) -lse=1 -mask=0 -num_splits=0 -kname=1 -v=1 bin/tile_example_fmha_fwd -prec=fp16 -mode=1 -page_block_size=128 -b=1 -h=32 -h_k=32 -d=128 -s=1 -s_k=$((75 * 128)) -lse=1 -mask=0 -num_splits=0 -kname=1 -v=1 bin/tile_example_fmha_fwd -prec=fp16 -mode=1 -page_block_size=128 -b=1 -h=8 -h_k=1 -d=256 -s=1 -s_k=$((3133 * 128)) -lse=1 -mask=0 -num_splits=0 -kname=1 -v=1 ``` ## Test Result All the tests must pass. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

[CK Tile] Fix V6 pipeline applicability and split-image initialization (#7936) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit ## Motivation After adding code generation via CK Tile Dispatcher, some fwd and bwd weight tests for CK Tile convolutions are failing. This PR introduced correct applicability checks and fixes the split-image parameter initialization such that non-applicable instances are not invoked during test execution and split-image instances are correctly initialized. ## Technical Details Investigation revealed two distinct problems 1. For bwd weight, the compute V3 uses prefetch of 3 distinct tiles, which works incorrectly when the number of K-slices addressed by the workgroup is 1. This occurs when a large split-K value is used for a problem that results in a small Gemm-K value. 2. For fwd direction, the current CK Profiler/test infrastructure doesn't initialize the split-image parameters for instance where split-image is enable. Uninitialized split-image values result in non-deterministic behavior where the tests might randomly fail. Fixed problem 1. by adding a check in `IsSupportedArgument` that marks the instance invalid if the `num_loops = ceil(GemmK / (k_batch * KPerBlock)) < 4` for V6 pipeline kernel instances. The check is compile-time eliminated for other kernels. Fixed problem 2. by adding initialization of split-image parameters when split-image is enabled. The default initialization corresponds to full image with no split, i.e., the number of splits is 1 and it has the size of the full image. Added unit tests for the added logic. ## Test Plan Running the following test suites cover the logic added in this PR - test_grouped_convnd_fwd_tile - test_ck_tile_grouped_conv_fwd - test_grouped_convnd_bwd_weight_tile - test_ck_tile_grouped_conv_bwd_weight All test suites above are included in the automated test runs. ## Test Result  ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

=?UTF-8?q?[CK=20TILE]=20Unification=20Work=20=E2=80=93=20?= =?UTF-8?q?More=20accurate=20tests=20for=20MmaPipelines=20(#6212)?= MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit ## Motivation This PR solves several issues: #### More accurate tests for MmaPipelines The current tests for the MmaPipelines (test_amdgcn_sparse_mma, test_amdgcn_wavewise_mma) use explicit input fragment vectors filled with 1s, and only check the output of a single lane. We should have tests that actually use the MmaPipelines with non-trivial input matrices and verify the complete output. Some other aspects of the current MmaPipelines tests that I noticed and deserve some attention: 1. There is sometimes iteration over K outside of the pipeline, which is then included in WaveTileK or FragK, which is not correct. We should remove it, move K iteration inside of the pipeline, or be more clear about this outer-K loop size and how it propagates downwards. 2. There is very tight coupling between the kernel, gtest code, and test_pipeline helper, requiring a lot of information and functions to be passed back and forth. 3. The test_pipeline helper is doing a bunch of register-related logic on the host (related to point 1) 4. Without this register logic the only thing it does is check the device, call the kernel, and check the output, but with a lot of boilerplate. #### Test helper for detecting target arch at HOST runtime There is a really apparent issue we faced while writing tests: Scenario: 1. Compile a test that supports both gfx950 and gfx1201 for gfx950 2. Run the test on a server that only has gfx1201 GPU Actual: Segmentation fault Expected: The test can correctly detect from HOST runtime that the DEVICE target_id was different and skips the test. Notes: The only way of detecting the COMPILER_TARGET_ID in the existing "arch" framework is launching a kernel and calling `get_compiler_target()` (so, from a DEVICE code). This will create a segmentation fault if the current arch differs from the target arch. To cope with this issue, we propose to export the compiler target(s) (note they can be many) through `projects/composablekernel/test/ck_tile/core/arch/CMakeLists.txt` and define a test helper to deal with such cases. #### Add composition support to Transforms We have a small number of Transforms which act on MmaOp input and output data, before and after the MmaOp call respectively. These are currently implemented to work on an MmaTile level, but in theory they are also supposed to work at a WaveTile level, i.e. after composition of multiple MmaTiles to create larger effective MNK dimensions. Currently the composed MmaTiles look like 2D C-style arrays of the individual MmaTile level register vectors (see WaveWiseMmaPipeline). The transforms should be able to take these and perform the proper transforms to the whole WaveTile at once. This might allow for better performing transformations. Note: This PR handles the SparseTransform case and if we don't end up doing scale as a transformation, there isn't really much left to do. If we end up having only the sparse transform as a non-trivial transform, then we could also consider removing the Transform framework.

[CK] Grouped conv profiler updates ## Motivation Reduce profiling time for no verification. ## Technical Details Remove not needed code for no verification ## Test Plan test_grouped_convnd* ## Test Result pending ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. AICK-1230

[GFX1250][CK_TILE] Add scale16 warp gemm unit tests ## Summary - Add scale16 WMMA intrinsic overloads and int64_t forwarding to warp gemm layers for gfx1250 - Add comprehensive wave-level unit tests for scale16 warp gemm (16x16x128 and 32x32x128 tile sizes) - Test all fp8/bf8 type combinations and TransposeC variants - Fix WarpGemm wrapper for non-uniform scale16 configurations Stacked on #7724 (FillUniformScaleDistribution / MX GEMM scale init). Pipeline enablement follows in the next PR.

[CK TILE][Windows] add `msvc::no_unique_address` support for Windows (#7786) ## Motivation While building Flash Attention 2 with CK backend, this warning will spam in every kernel: ``` DEBUG [1/1837] hipcc.exe ... DEBUG In file included from H:\ROCm\flash-attention\build\fmha_fwd_d32_bf16_batch_b64x64x16x32x32x32_r4x1x1_r4x1x1_w16x16x16_w16x16x16_qr_vr_pssk_nlogits_alibi_mask_lse_ndropout_nskip_nqscale_ntrload_nsink_gfx12.cu:6: DEBUG In file included from H:\ROCm\flash-attention\csrc\composable_kernel\example\ck_tile\01_fmha\fmha_fwd.hpp:6: DEBUG In file included from H:\ROCm\flash-attention\csrc\composable_kernel\include\ck_tile/core.hpp:111: DEBUG H:\ROCm\flash-attention\csrc\composable_kernel\include\ck_tile/core/tensor/tile_scatter_gather.hpp:1246:7: warning: unknown attribute 'no_unique_address' ignored [-Wunknown-attributes] DEBUG 1246 | [[no_unique_address]] std::conditional_t<kUseGlobalLoad_, PageIdxArray, gl_field_empty_t> DEBUG | ^~~~~~~~~~~~~~~~~ DEBUG H:\ROCm\flash-attention\csrc\composable_kernel\include\ck_tile/core/tensor/tile_scatter_gather.hpp:1254:7: warning: unknown attribute 'no_unique_address' ignored [-Wunknown-attributes] DEBUG 1254 | [[no_unique_address]] std::conditional_t<kUseGlobalLoad_, index_t, gl_field_empty_t> DEBUG | ^~~~~~~~~~~~~~~~~ DEBUG 2 warnings generated when compiling for host. ... ``` ## Technical Details `[[no_unique_address]]` is not working on Windows LLVM, should use `[[msvc::no_unique_address]]`. ## Test Plan Build FA2 with CK backend. ## Test Result No warnings, no errors. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. Co-authored-by: Illia Silin <98187287+illsilin@users.noreply.github.com>

composablekernel: remove stray *.hpp.bk backup artifacts (#7974) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Four `*.hpp.bk` files were accidentally committed to `projects/composablekernel/`, likely as leftovers from a prior merge or conflict resolution. Each is an older snapshot of its `.hpp` counterpart — the canonical `.hpp` files are newer and contain the correct current content. ## Deleted files | File | vs. `.hpp` counterpart | |---|---| | `ck_tile/core/tensor/tile_window.hpp.bk` | Older version: uses legacy `bool isL1Cache`/`PrefetchL1` template params; missing `DataCachePrefetchKind`-based prefetch API and `data_cache_prefetch.hpp` include | | `ck_tile/core/tensor/load_tile_transpose.hpp.bk` | Older version: missing `#if defined(__gfx950__)` guard and `Quad` struct (~90 lines) for gfx1250 architecture | | `ck_tile/ops/gemm/warp/warp_gemm_dispatcher.hpp.bk` | Older version: missing `WmmaTag`, `IsScale16` template param, and several newer dispatcher specializations | | `ck_tile/ops/gemm_quant/block/block_universal_gemm_as_bs_bquant_cr.hpp.bk` | Older version: `KPackA`/`KPackB` (since renamed `KPack`); uses `static_ford` (since refactored to nested `static_for`) | ## Verification - No other `.bk` files exist in `projects/composablekernel/`. - No build scripts, CMake files, includes, or documentation reference these `.bk` files. - No `.hpp` files were modified.

[ck] Enforce ASCII-only C/C++ sources for hipRTC compatibility (#7829) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit ## Summary CK source files must be compilable via **hipRTC (HIP runtime compilation)**, whose preprocessor does not accept non-ASCII bytes anywhere in a translation unit — **including in comments**. Bytes that are harmless under `hipcc` (em-dashes, smart quotes, multiplication signs, Greek letters, box-drawing glyphs, etc.) cause hipRTC to fail at preprocessing time. These regularly leak in via LLM-assisted authoring or copy/paste from formatted documents and silently break hipRTC paths that are not exercised by the default `hipcc`-based build matrix. This PR (a) cleans every existing violation (53 files) and (b) adds a pre-checkin gate so new violations are rejected before merge. ## File extensions covered Both the cleanup scan and the new Jenkins enforcement stage use the same predicate: ``` *.h *.hpp *.cpp *.h.in *.hpp.in *.cpp.in *.inc *.cl ``` (excluding `*/build/*` and `*/include/rapidjson/*`). This is a strict superset of the existing `Clang Format` stage's predicate — `*.inc` is added so test-fixture include files are also gated. The local pre-commit hook's `c++/inc` type filter covers the same set. ## Why no enforcement today CK is opted out of the rocm-libraries root `.pre-commit-config.yaml`, so the existing `pre-commit` workflow doesn't touch CK. The local CK `.pre-commit-config.yaml` only runs for developers who installed hooks. The **authoritative gate is therefore the new Jenkins stage** in this PR; the local hook is convenience. ## Commit layout (bisect-friendly) 1. `79798aa6261` — **`[ck] Convert reflect/ rendering to ASCII for hipRTC compatibility`** Behavior change, isolated. `TreeFormatter` swaps `├─ / └─ / │ ` for `|- / +- / | ` (3-col width preserved so alignment is unchanged). `conv_description.hpp` swaps `×` for `x` as the dimension separator. `test_conv_description.cpp` expected strings updated in lockstep so the snapshot test stays green. This is the only commit in the series with observable runtime impact. 2. `738fdb0d81c` — **`[ck] Strip non-ASCII bytes from C++ sources for hipRTC compatibility`** Mechanical text cleanup across 53 files. Replacements happen in comments or in `std::cout` strings that are not asserted on by any test. None of the 174 `.inc` files in the tree required edits, but they were in the scan's predicate so the enforcement stage's predicate is a superset of what was scanned. Full replacement table in the commit message. 3. `1d7cd8ba235` — **`[ck] Enforce ASCII-only C/C++ sources for hipRTC compatibility`** - New `projects/composablekernel/script/check_ascii_only.sh` (modeled on `check_copyright_year.sh`). - New entry in `projects/composablekernel/.pre-commit-config.yaml` under the local-hooks block (`types_or: [c++, inc]`). - New `ASCII Only Check` parallel stage in `projects/composablekernel/Jenkinsfile`'s `Static checks` block, mirroring the existing `Clang Format` stage but with `*.inc` added to the find predicate. Always-on, no `RUN_CPPCHECK` gate. The tree is buildable at every commit boundary. Commit 1 leaves 50 known violations; commit 2 leaves 0; commit 3 wires the gate. ## Demo Script output on a synthesized violation: ``` $ printf '// em-dash test \xe2\x80\x94 here\n' > /tmp/bad.cpp $ projects/composablekernel/script/check_ascii_only.sh /tmp/bad.cpp ERROR: /tmp/bad.cpp contains non-ASCII bytes: 1:// em-dash test — here Fix: replace with ASCII (em-dash -> --, smart quotes -> ", arrows -> ->, etc.) $ echo $? 1 ``` Full repo scan after the cleanup commits (note the `-name '*.inc'` clause): ``` $ cd projects/composablekernel && find . -type f $ -name '*.h' -o -name '*.hpp' -o -name '*.cpp' \ -o -name '*.h.in' -o -name '*.hpp.in' -o -name '*.cpp.in' -o -name '*.inc' -o -name '*.cl' $ \ -not -path '*/build/*' -not -path '*/include/rapidjson/*' -print0 \ | xargs -0 -P 8 -n 64 script/check_ascii_only.sh $ echo $? 0 ``` ## Test plan - [ ] Jenkins PR build: confirm new `Static checks -> ASCII Only Check` stage runs green over the full predicate (incl. `*.inc`) and existing `Clang Format` stage is unaffected. - [ ] `test_conv_description` passes against the ASCII tree-formatter output (touched in commit 1). - [ ] Local: `pre-commit run ascii-only-checker --all-files` runs cleanly after installing CK pre-commit hooks via `script/install_precommit.sh`. - [ ] Manually inject a non-ASCII byte in any `.cpp/.hpp/.inc` file, push: confirm Jenkins fails the new stage with a clear error. - [ ] Spot-check a representative subset of touched files under hipRTC compilation to confirm no remaining hipRTC-blocking content (optional, since the static byte check is a sufficient condition for hipRTC preprocessor acceptance on this dimension). 🤖 Generated with [Claude Code](https://claude.com/claude-code)

[CK] Fix latest build issues with staging compiler. ## Motivation Fixing new warnings with staging compiler. ## Technical Details  ## Test Plan  ## Test Result  ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

[CK] Load ck.groovy via Jenkins Shared Library MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit ## Motivation This allows the CI service to have a configuration source-of-truth outside the PR under test, allowing rapid system changes. Bug fixes on the develop branch propagate immediately to all pipelines that don't override the parameter -- no rebase required. A new `USE_CURRENT_BRANCH_FOR_CK_GROOVY` parameter lets contributors test pipeline changes on their own branch without any extra configuration. ## Technical Details - `loadCk()` in the Jenkinsfile is updated to call `library("ck@${branch}").ck.get()` instead of `checkout scm` + `load "vars/ck.groovy"`. The `checkout scm` inside `loadCk()` is removed since Jenkins now handles the library fetch internally. - A `USE_CURRENT_BRANCH_FOR_CK_GROOVY` boolean parameter (default: off) is added. When off, `ck.groovy` is always loaded from `develop` — all normal PR builds are unaffected. When on, `ck.groovy` is loaded from the current branch automatically via `env.CHANGE_BRANCH`, so contributors testing pipeline changes just tick the box. - `return this` is removed from the end of `ck.groovy`. This was required by the `load` convention but is not needed (and can cause errors) in a shared library context. - `loadCk()` is kept at every call site rather than called once at the top, preserving restart-from-stage safety — if a build is restarted from a mid-pipeline stage, `ck` is still initialized correctly. - The Jenkins Shared Library named `"ck"` must be registered in Jenkins Global Pipeline Libraries ## Test Plan 1. Trigger "Build with Parameters" on the PR branch with `USE_CURRENT_BRANCH_FOR_CK_GROOVY=true` 2. Verify "Determine CI Execution" stage completes and the library() calls indicates the current branch 3. Verify "Static checks" stage completes. 4. Trigger a second build with `USE_CURRENT_BRANCH_FOR_CK_GROOVY=false` (default) to confirm normal builds still load from `develop`. ## Test Result Verified both paths. The develop library is loaded by default, the branch library is loaded when the parameter is enabled. ## Submission Checklist - [ X ] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

[ck] Updated CK Tile documentation to use mermaid diagrams (#7955) ## Motivation There were mermaid diagrams in the CK Tile doc that were converted to svg. However, there is an extension for mermaid diagrams. The conf.py and requirements.in have been updated to use that extension instead of the svg files.

[CK Tile] PermuteN support MX GEMM ## Motivation Add PermuteN support to preshuffle MX GEMM ## Technical Details - Modify `shuffle_b_permuteN` to support MX preshuffled layout - Add `preShuffleScalePermuteN` with same functionality of `preShuffleScale` but layout consistent with PermuteN - Include MX pre-processing functions in the library ## Test Plan Add test configuration for permuteN with preshuffle (both FP4 and FP8) ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. Co-authored-by: Cong Ma <congma13@amd.com>

[CK_Tile] Add wmma_bf16f32_16x16x32_bf16 via fused-downconvert override (#8028) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit ## Summary Adds `__builtin_amdgcn_wmma_bf16f32_16x16x32_bf16` (fp32 accumulate → bf16 output) to the CK Tile WMMA warp-gemm path. **API only** — the unit test is split into a stacked PR (#8035) so this API change can be reviewed in isolation. ## Changes (4 files) - **16-bit trait:** `wmma_intrinsic_downconvert` (calls the bf16f32 builtin — fp32 C in, bf16 C out) plus `COutDataType = bf16_t` / `COutVecType`. - **`WarpGemmAttributeWmmaImpl` / `WarpGemmAttributeWmma`:** `mac_downconvert(c_fp32, a, b)` (kTransC-aware) returning the bf16 C-output vector. - **`WarpGemmImpl`:** `mac_downconvert` tail handler producing a bf16 C-output tile from the fp32 accumulator tile, reusing `CWarpDstrEncoding` (output layout identical to the f32 C tile). Verified on gfx1250 (via the stacked test PR #8035): the test passes; the existing WMMA warp-gemm test is unaffected (additive change only).

[CK Tile] Async support preshuffle GEMM ## Motivation Add async support to existing preshuffle GEMM pipeline ## Technical Details Notes: the implementation avoids previous strategy of duplicating pipelines for async support and instead add a switch `Async` to the ops Problem to enable async pipeline. Then, integrate the async pipeline in the existing one. This allows to avoid code duplication and facilitate the integration of buffer load to lds in existing pipelines. In my opinion, it should be used also for other pipelines which don't support buffer load to lds yet and it would also be a good idea to refactor the existing async GEMM pipelines with the same approach. Summary: - integrate buffer load to lds in existing pipeline - add optimal tensor descriptors for vmem loading and lds reading. They are currently optimized for 16x16 wave tiles but they also work for 32x32 wave tiles. Optimizations for 32x32 wave tile requires different lds layout and it will be done in a follow-up issue - Add async config to examples - Add test (gfx950 only) ## Test Plan New test for gfx950 `test_ck_tile_gemm_pipeline_wp_async` ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

[CK_TILE] [QuantGEMM] Fix SplitK tail handling and other improvements (#7199) This pull request introduces improved and more robust split-K support for quantized GEMM. The main changes add runtime validation, utility functions for split-K batch calculations, pointer offset handling for split-K in grouped kernels, and enhanced support for various tensor layouts. The changes also improve error handling and provide more flexibility for runtime tail handling in split-K pipelines. **Split-K Support and Validation Enhancements:** * Added runtime validation to ensure `k_batch` is a positive integer and that split-K configurations do not produce empty final batches or mismatched pipeline tails, with detailed error messages and logging for misconfiguration. [[1]](diffhunk://#diff-d000149a681cd42bfb9947872c603e556cea26cbd7fd4f8f60afc6595d975871R1184-R1211) [[2]](diffhunk://#diff-d000149a681cd42bfb9947872c603e556cea26cbd7fd4f8f60afc6595d975871L1161-R1250) * Introduced utility functions `get_splitk_batch_k_read` and `get_splitk_last_batch_k` to compute per-batch K read sizes and handle split rounding, ensuring correct and consistent split-K batch partitioning. [[1]](diffhunk://#diff-d000149a681cd42bfb9947872c603e556cea26cbd7fd4f8f60afc6595d975871R206-R234) [[2]](diffhunk://#diff-635b89bdffa96b2b42f1632520cde36701d7d631e864185591f6b32f7645cf47L104-R107) [[3]](diffhunk://#diff-d000149a681cd42bfb9947872c603e556cea26cbd7fd4f8f60afc6595d975871L388-R417) [[4]](diffhunk://#diff-d000149a681cd42bfb9947872c603e556cea26cbd7fd4f8f60afc6595d975871L1161-R1250) * Changed the default value of `k_batch` in `QuantGemmHostArgs` to 1 (no split-K) for safer default behavior. **Pointer Offsets and Grouped Kernel Handling:** * Updated `QuantGroupedGemmKernel` to apply split-K per-batch offsets to all input pointers, mirroring the behavior of non-grouped kernels and ensuring correctness for split-K launches. * Modified AQ tensor view handling to correctly reflect the remaining K-groups from the split-K batch's offset position, improving accuracy for split-K in grouped kernels. **Pipeline and Layout Flexibility:** * Added support for runtime selection of split-K tail handling via a new template parameter `RuntimeSplitKTail_`, with new helper methods to dispatch GEMM pipelines accordingly. [[1]](diffhunk://#diff-d000149a681cd42bfb9947872c603e556cea26cbd7fd4f8f60afc6595d975871R273) [[2]](diffhunk://#diff-d000149a681cd42bfb9947872c603e556cea26cbd7fd4f8f60afc6595d975871R1496-R1567) [[3]](diffhunk://#diff-d000149a681cd42bfb9947872c603e556cea26cbd7fd4f8f60afc6595d975871L1427) [[4]](diffhunk://#diff-d000149a681cd42bfb9947872c603e556cea26cbd7fd4f8f60afc6595d975871L1447-R1629) [[5]](diffhunk://#diff-d000149a681cd42bfb9947872c603e556cea26cbd7fd4f8f60afc6595d975871L1459-R1641) * Improved handling for tensor layout cases, including preshuffled B and both row-major and column-major AQ layouts, ensuring correct pointer arithmetic and compatibility checks. [[1]](diffhunk://#diff-d000149a681cd42bfb9947872c603e556cea26cbd7fd4f8f60afc6595d975871R438-R454) [[2]](diffhunk://#diff-d000149a681cd42bfb9947872c603e556cea26cbd7fd4f8f60afc6595d975871L464-R516) [[3]](diffhunk://#diff-d000149a681cd42bfb9947872c603e556cea26cbd7fd4f8f60afc6595d975871R1184-R1211)

fix(ck): Stream-K Tile Engine GPU Query Fix ## Motivation The Stream-K tile engine validation utilities call rocminfo at build time to detect the GPU architecture, which prevents building on CPU-only nodes. The GPU target is already known from CMake's SUPPORTED_GPU_TARGETS, so runtime hardware detection is unnecessary during code generation. ## Technical Details The gemm_streamk_validation_utils.py file used subprocess.check_output(["rocminfo"]) to query GPU hardware at CMake configure time. This call originated from get_gpu_name_by_id() and used during tile configuration validation. On CPU-only build nodes, this fails because rocminfo either doesn't exist or returns no GPU devices. The main changes are as follows: - Removed runtime GPU detection infrastructure: deleted get_gpu_name_by_id(), set_gpu_targets(), get_configured_gpu_targets(), the _configured_gpu_targets module variable, and the GPU_NAME_PATTERN regex. - Added gpu_target as an explicit parameter. is_tile_config_valid(), validate_gemm(), validate_warp_tile_combination(), validate_warp_configuration(), and validate_lds_capacity() now accept gpu_target as a required parameter instead of querying hardware internally. The corresponding changes were also made in gemm_streamk_instance_builder.py and CMakeLists.txt - Added WARP_SUPPORTED_COMBINATIONS for per-GPU warp config validation, and LDS_SIZE_MAP / DEFAULT_LDS_SIZE for GPU-aware LDS capacity checks. ## Test Plan The benchmarks were compiled and run on a GPU as well as a CPU only node to verify correctness. ## Test Result All tests passed ## Related JIRA ID : AICK-1635 ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

fix(ck): correct GU-fusion B-up vector loading ## Motivation Updating latest CK in Aiter will cause ATOM Qwen tests failure. Bisect CK commit narrow down to the changes in ROCm/rocm-libraries#4798. JIRA ID AICK-1709 ## Technical Details The fix is to pass b_thread_vec_up instead of b_thread_vec in thread_buf_to_vec_loader for LoadBUp. ## Test Plan default CK CI test manual testing for ATOM Qwen tests. ## Test Result default CK CI test pass ATOM Qwen tests pass ## Submission Checklist - [ ] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

Users/mkulikow/ck/data prefetch in mxgemm pipeline JIRA ID : AICK-1670 ## Motivation Added example for data cache prefetch in mx gemm pipeline while also fixing some bugs in data cache prefetch pipeline ## Technical Details Add a standalone flatmm example (mx_flatmm_data_cache_prefetch) that runs MX GEMM through the compute TDM v1 pipeline (GemmPipelineAgBgCrCompTDMV1) with hardware data cache prefetch on gfx1250. Prefetch destination is selectable per operand (A/B) between L1, L2 or None via the DataCachePrefetchKind trait, exposed through -prefetch_a_l1 / -prefetch_b_l1 CLI flags, with an optional -compare mode against a no-prefetch run. Guarded behind gfx125 in the 18_flatmm CMakeLists. Also fix data cache prefetch being silently disabled in the scaled operator() of GemmPipelineAgBgCrCompTDMV1. The scaled path defaulted data_cache_prefetch_a/b to false and only set them under UseClusterLaunch, so with cluster launch off the runtime guards folded away every prefetch and no global_prefetch_b8 was emitted despite an L1/L2 policy. Default them to true (matching the non-scaled operator()); emission stays gated by the compile-time UseDataCachePrefetch policy check, so None still emits nothing. ## Test Plan Checked on simulators: test name: tile_example_mx_flatmm_mxgemm_data_cache_prefetch ## Test Result fp4 for 512x512x4096: no prefetch/L1 prefetch: 71,334 / 54,889 ( 23.1% increase ) ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests.

fix(ck-tile): Prefer using amd-smi over rocm-smi ## Motivation CK tooling invokes `rocm-smi` directly in many places for GPU discovery and info. `amd-smi` is the modern replacement and is preferred going forward, but call sites should not need to know which tool is present. This PR centralizes GPU SMI access behind a single set of wrappers that prefer `amd-smi` and transparently fall back to `rocm-smi`, so every call site gets consistent behavior with no duplicated detection logic. ## Technical Details - Added `tile_engine/ops/common/smi_utils.py` as the single source of truth for SMI access: parsers for both tools plus wrappers `detect_gpu_ids()`, `count_gpus()`, `show_gpu_info()`, `check_gpu_available()`, and `show_version()`. - Tool order is chosen by `_smi_order()`: `amd-smi` first, `rocm-smi` fallback if `amd-smi` is missing or fails. `CK_SMI_TOOL=rocm-smi` forces rocm-smi first (mainly for testing). - Added `tile_engine/ops/common/smi_cli.py`, an argparse CLI (`list-ids`, `count`, `show-info`, `check`, `show-version`) so shell scripts can reuse the same Python logic. - Added thin `ck_smi_*` delegates in `script/tools/common.sh` that call `smi_cli.py` (no parsing logic in Bash). - Migrated call sites off direct `rocm-smi`: `gemm_full_benchmark.py` now uses `detect_gpu_ids()`; `generate_test_dataset.sh`, `ck-status`, `ck-start`, `ck-docker`, `ck-exec`, `ck-shell`, `ck-rocprof.md`, and the GEMM `README.md` now use the wrappers. ## Test Plan - `python3 -m unittest tile_engine.ops.common.test_smi_utils -v` live tests comparing `rocm-smi` vs `amd-smi` normalized fields (GPU IDs, product, gfx, driver) and verifying each wrapper against live output. - `CK_SMI_TOOL=rocm-smi python3 -m unittest tile_engine.ops.common.test_smi_utils -v` verifies the rocm-smi override path. - `bash script/tools/test_ck_smi_helpers.sh` pure-bash live comparison of the two tools' fields (no Python dependency). - Ran on a GPU host. ## Test Result All 11 unittest cases pass; the bash comparison reports all fields matching. Both the default (amd-smi first) and `CK_SMI_TOOL=rocm-smi` paths return identical GPU IDs. ## Submission Checklist - [x] Look over the contributing guidelines at https://github.com/ROCm/ROCm/blob/develop/CONTRIBUTING.md#pull-requests. JIRA ID : AICK-1649

fix(ck-tile): repair #9308 merge truncations in gemm_utils + unified_gemm_codegen (develop broken) (#10105) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit ISSUE ID: #9308 Fixes the develop breakage introduced by the #9308 multi-D merge (commit 7fcb5f36). When I was fixing merge conflict, some line was deleted and CI did not test it ( bridge test was not added CI that time). These two issue has been fixed with this PR. ## Summary The multi-D merge **#9308** (commit `7fcb5f36`) left **two independent truncations** on `develop`, both of which make the dispatcher Python fail to parse: 1. `dispatcher/python/gemm_utils.py` — `import gemm_utils` raises `SyntaxError: '(' was never closed`. 2. `dispatcher/codegen/unified_gemm_codegen.py` — `SyntaxError: unterminated string literal` (multi_d / multi_abd codegen can't run). Both verified against the GitHub `develop` blob. Every consumer of the multi-D / multi-ABD Python paths is broken today. ## Root cause - **gemm_utils.py:** the multi-D `run()` landed **inside `GpuMultiABDRunner`** and was truncated at an unterminated `return MultiDGemmResult(`; `GpuMultiDGemmRunner` was left with only `__init__`. - **unified_gemm_codegen.py:** `_multi_d_single_include()` returns an f-string of C++ `#define`/exports whose closing `"""` was dropped. ## Why CI stayed green The dispatcher Python tests that import these modules were **not registered in `dispatcher/tests/CMakeLists.txt`**, so they never ran in the gate. ## Changes - **gemm_utils.py:** move multi-D `run()` (+ `kernel_name`/`num_d_tensors`) into `GpuMultiDGemmRunner` and complete the `MultiDGemmResult(...)` return; remove the stray block from `GpuMultiABDRunner`. - **unified_gemm_codegen.py:** close the truncated f-string in `_multi_d_single_include`. - **Tests + CI:** add `TestModuleImportsAndRunnerShape` to `test_gemm_utils.py` (import canary + multi-D runner-shape assertions + `ast.parse` canary over `unified_gemm_codegen.py`), and **register `test_gemm_utils.py` in CMake (`dispatcher_test_gemm_utils`)** so ctest/CI runs it. ## Test plan - [x] `import gemm_utils` succeeds (was SyntaxError) - [x] `ast.parse` of `unified_gemm_codegen.py` succeeds (was SyntaxError) - [x] `python3 -m unittest discover -p test_gemm_utils.py` -> 21 passed - [ ] CI runs `dispatcher_test_gemm_utils` (newly wired) ## Follow-up (separate) Enable the dispatcher Python test suite broadly in the PR CI gate to fully close the coverage gap.

…dormant, Stage 27) All new paths gated behind GGML_HIP_* env vars -> OFF by default. Zero change to default dispatch. Backs the AMD PR ROCm/composable_kernel#3759 §3-4 journey. WINS: - mul_mat_2of4_fp8_mmq: MMQ-grade 2:4-sparse fp8 GEMM (cooperative tiles, LDS staging, ILP>=4). At/ahead of our native-fp8 WMMA MMQ on the 3 dominant GEMM shapes (+1.8% gate-up / +3.9% q-o / +32% down-proj). - mul_mat_dense_fp8_mmq: dense-fp8 MMQ twin. - shuffle-based activation-quantize in mmvq.cu (128-thread, float4 loads, warp-shuffle max-reduction, zero LDS/__syncthreads) -- replaces the 32-thread LDS+sync reducer; ~36us->18us, closed ~40% of the remaining 8% gap. Both kernels now high-90s% of native-fp8 (0.989x same-session). - k/v adaptive tile in mul_mat_2of4_fp8.cu. CORRECTIONS / MEASUREMENT CONTROLS (the honest half): - mul_mat_dense_fp8_v3: dense-fp8 twin used as the sparsity-isolation control. Proves the ISA 2x/4x sparsity ceiling WASHES to ~1% at library grade (2x ISA -> 1.25x crude -> ~1.01x well-tiled). This is the finding sent to AMD. - swmmac24_iu4_fixed: NEGATIVE test. Applying CK's idx swap+XOR-1 to our OWN correct native SWMMAC encoder BREAKS it (err 0 -> 385). Proves the CK bug is compress/packing-convention-specific, NOT universal hardware behavior -- folded into PR #3759 Bug 3 precision note. - int4_24_probe: int4-2:4 probe (FAILING, gated off). Source of the "4x" decomposition: 2x(int4 packing, already in dense wmma_i32_16x16x32_iu4) x 2x(2:4 sparsity, washes) -> net ~1.2x weight-bandwidth only.

doplxyz · 2026-08-14T23:29:54Z

Thanks for digging this out and writing it up in this much detail — this had been sitting without a
review for three weeks, and the analysis in the description made it possible to check the claims
rather than guess at them. I have a gfx1201 box, so I ran it.

Short version: bug 1's fix holds up — I reproduced a failure on the base tree and, for what this
test measures, traced it to that one line. But the committed test can't be built from the PR alone, it passes on
the unfixed tree in the configuration that is buildable, and it never reaches the iu4 path that
bugs 2 and 3 are about. Separately, the changed code is not restricted to gfx1201, and I think that's
the thing to sort out before merge.

This is not a formal approval: I can only speak for gfx1201 and the int8 path.

Environment. AMD Radeon RX 9070 XT, gfx1201, amdgcn-amd-amdhsa--gfx1201; Linux 6.14.0-37,
amdgpu 6.19.14.31400100; container
rocm/pytorch:rocm7.14_ubuntu24.04_py3.12_pytorch_release_2.12.0, AMD clang 23.0.0git
(ROCm/llvm-project 46fcb339fb61). Trees compared: base
8fc1ac24e9bd7b431663a15e2122ce02c2979d37 — which is git merge-base of this PR's head and
develop, and base..head is exactly the three files in this PR — versus head
ac24ac28d662acf279478545cec541d4dde00f31. Identical test source and compile flags on both sides,
separate clean trees and build directories, three fresh processes per variant, identical stdout on
every repeat.

Two environment-side notes so the commands reproduce, neither of which is a problem with this PR:
that container needs --rocm-device-lib-path=<sdk>/lib/llvm/amdgcn/bitcode, and it ships only
libamdhip64.so.7, so a libamdhip64.so symlink is needed for the link step.

1. The repro can't be built from the PR alone

sparse_swmmac_correctness_repro.cpp:34 has

#include "real_a_tile.h"  // REAL_A_16x128: Quark int8 2:4 weight tile

outside the #ifdef USE_REAL_TILE at line 264, and real_a_tile.h isn't one of the PR's three files.
From a clean checkout the translation unit doesn't compile either way, so the max_abs_err=352
figures can't be checked by a reviewer.

I generated a substitute tile to get something running: same int8_t[16][128] shape, at most two
non-zeros per group of four along K, covering all six two-survivor position pairs, all four
single-survivor positions and the zero-survivor case, with values derived from (row, group, position) so that a misplaced survivor is more likely to show up numerically rather than cancel.
This is my input, not yours — nothing below should be compared against 352.

2. base fails, head passes

variant	K=32	K=64	K=128
base `8fc1ac2`	`max_abs_err=56`	`76`	`92`	FAIL
head `ac24ac2`	`0`	`0`	`0`	PASS

The harness wraps hipMalloc / hipMemcpy / hipMemset / hipDeviceSynchronize / hipFree in
HIP_CHECK_ERROR, and none tripped, so the failing numbers are a real kernel result rather than a
silent launch failure.

3. The whole effect is bug 1

Two more variants — base plus only the nonzero_elems true-zero default, and head with only that
line reverted:

variant	K=32	K=64	K=128
base + bug-1 line only	`0`	`0`	`0`	PASS
head − bug-1 line	`56`	`76`	`92`	FAIL

The head − bug-1 stdout is identical to base's, byte for byte (diff on the run logs). So for the
failure this test measures, on these two commits, with this input and these three shapes, that one
line is both necessary and sufficient. I'm not claiming more than that — this doesn't establish
correctness over all valid inputs, types or architectures. The reasoning in your comment matches what
I see.

4. The test never reaches iu4, so bugs 2 and 3 get no numerical coverage from it

I disassembled the code objects of the kernels that actually launch, rather than grepping the
executable. The three SparseGemmKernel symbols (WaveTileK = 32 / 64 / 128) contain 1, 2 and 4
SWMMAC instructions respectively, zero v_wmma, zero v_mfma — and all seven are
v_swmmac_i32_16x16x32_iu8. There is no iu4 instruction anywhere in the binary.

That follows from the instantiation: the test uses SparseMmaPipeline<int8_t, int8_t, int32_t, ...>,
and int8_t picks up the generic numeric_traits::PackedSize == 1, so the if constexpr(PackedSize == 1) branch is taken, the packed-nibble path of bug 2 is compiled out, the SWAP + XOR-1 transform of
bug 3 is never reached, and TotalCompressedElems * MmaOp::APackedSize evaluates unchanged.

So the description reasons about all three bugs from the iu4 side, but the committed test only
exercises iu8. I haven't verified bugs 2 or 3 either — building an iu4 oracle independently of the
transform under test (nibble order, sign extension, logical vs packed K) is its own job, and reusing
your transform as the oracle would be circular.

What I could check is the shape side, with a probe that instantiates
SparseMmaPipeline<pk_int4_t, pk_int4_t, int32_t, ...> directly and launches it:

	base	head
`TotalCompressedElems`, K=32	4	8
`TotalCompressedElems`, K=64	8	16
`TotalUncompressedElems`	8 / 16	unchanged
`IdxNumWords`	1	unchanged at these shapes

Both trees compile, both emit v_swmmac_i32_16x16x32_iu4 and v_swmmac_i32_16x16x64_iu4, and
hipDeviceSynchronize() returns success. That's consistent with your ISA reading — 8 index values
per lane for a K=32 iu4 tile — but it's a shape and codegen check, not a correctness one.

5. The buildable configuration passes on the unfixed tree — and the reason is the one you identified

Built without -DUSE_REAL_TILE, so the #else synthetic path runs. (The substitute header from
§1 is still needed even here, since the #include is unconditional — dropping the define alone does
not make the file compile.)

	result
base `8fc1ac2`, no `USE_REAL_TILE`	ALL PASS
head `ac24ac2`, no `USE_REAL_TILE`	ALL PASS

I expected that to be because the synthetic input has no under-filled groups, but that isn't it. I
replayed the exact host-side fill (mt19937(42), uniform_int_distribution(-8, 8), then
apply_sparse_pattern) and counted:

K	groups	0 non-zeros	1 non-zero	2 non-zeros
32	128	0	16	112
64	256	2	26	228
128	512	2	55	455

So under-filled groups are plentiful — the distribution includes 0 — and it still passes on the buggy
code. The reason is exactly the precondition you call out in the comment: apply_sparse_pattern
always zeroes slots 1 and 3, so survivors only ever sit at slots 0 and 2, and the old
{a_vec[i*4+2], a_vec[i*4+3]} default therefore always seeds slot 1 from a guaranteed zero. My tile
breaks it because survivors also sit at positions 2 and 3.

The practical consequence: once the missing header is supplied in whatever form, the configuration
that does not define USE_REAL_TILE reports a green run on the bug. Worth committing an input (or
generating one in the test) that puts survivors at position 3, and making that the default.

6. The change isn't scoped to gfx1201, and bug 1's fix demonstrably reaches CDNA

This is my main question before merge, and I don't think it's answerable from gfx1201 alone.

Neither changed hunk has an architecture predicate; both branch on PackedSize / APackedSize.
More to the point, the transforms selector in sparse_transforms.hpp:381 is specialized purely on
the op family:

struct MmaTransformsDefaultSelector<MmaOp, CompilerTarget,
                                    std::enable_if_t<MmaOp::OpFamily == MmaOpFamily::SPARSE>>
{ using SelectedTransforms = MmaDefaultTransformsSparse<MmaOp::kCompressionRatio>; };

with no enable_if_target_family_gfx*, unlike the dense gfx9 / gfx11 / gfx12 selectors right
alongside it. So every SPARSE op on every target resolves to this compress_a_impl. And
sparse/mfma/sparse_gfx9.hpp defines SPARSE ops for GFX942 / GFX950 over fp16_t, bf16_t,
int8_t, fp8_t, bf8_t — none of which specialize PackedSize, so they all take the same
PackedSize == 1 branch that bug 1's fix changes.

I checked that rather than inferring it. A compile-only probe that instantiates
SparseMmaPipeline<int8_t, int8_t, int32_t, 16, 16, 64, ..., Gfx942Target> and asserts

static_assert(std::is_same_v<
    typename MmaTransformsDefaultSelector<MmaOp942, Gfx942Target>::SelectedTransforms,
    MmaDefaultTransformsSparse<MmaOp942::kCompressionRatio>>);

compiles on both trees at --offload-arch=gfx942, and the resulting code object contains
v_smfmac_i32_16x16x64_i8. So that CDNA instantiation does route through the compress_a_impl this
PR changes, and takes the PackedSize == 1 branch.

That makes bug 1's fix a behaviour change for gfx942 as well, for inputs whose survivors don't sit
where the old default assumed. I think it's the correct fix and my gfx1201 result supports it — but
whether CDNA results actually change in practice, and whether anything downstream depended on the old
behaviour, is a runtime question I can't answer without the hardware.

The narrower half of this: I found no pk_int4_t SPARSE op for gfx9 in the tree, so bug 2's packed
branch and bug 3's SWAP + XOR-1 look gfx12-only in practice today, which limits the blast radius of
the empirically-derived metadata mapping considerably. It's still worth saying out loud that those
branches carry no architecture condition, so a future CDNA pk_int4 sparse op would silently inherit a
mapping that was measured on gfx1201.

On CI: the GitHub API reports no check suites and no commit statuses for either cd34d2f or
ac24ac2. I can't tell from outside whether anything ran elsewhere, but from the PR there's no
visible regression signal for the architectures the selector reaches. That's a project-side
infrastructure question rather than something to put on you.

7. Minor, non-blocking

The comments carry internal process labels (Stage-17b fix (local, not upstreamed),
Stage-17c CLOSURE fix, this Stage-17c audit's own probe) — worth rewording for upstream.
The new static_assert(PackedSize == 2 && is_same_v<LogicalADataType, pk_int4_t>) narrows the
packed path to pk_int4_t explicitly. I haven't tested whether any other packed A type reaches
this code today, so I don't know if it changes anything in practice — just noting it's a scope
change beyond the three bugs.
You already note the gtest port is pending. For what it's worth, test/ck_tile/CMakeLists.txt
enumerates its subdirectories with explicit add_subdirectory(...) calls rather than a glob, and
gfx1201_sparse_swmmac isn't among them, so the port will need that line too.

What would let me say more

Commit real_a_tile.h, or generate an equivalent tile inside the test, so the fail→pass is
reproducible from the PR alone — and make sure the default-built configuration is one that fails
before the fix.
Add an iu4 / pk_int4_t case, since bugs 2 and 3 have no numerical coverage without one.
Say what architecture scope is intended, given §6.

Happy to re-run any of this on gfx1201. If it would help, I can also try to pin bug 3's mapping
empirically and independently of sparse_transforms — a minimal wave-level kernel issuing raw
v_swmmac_*_iu4 with one-hot operands, enumerating the metadata encodings and reading the
contributing positions back out of the result — and then compare what that gives against SWAP + XOR-1.
That would be a second measurement rather than a derivation from the doc, which as you note doesn't
capture the iu4-specific detail. Say the word and I'll put the generator, flags, logs and per-symbol
disassembly somewhere you can pull them.

Address review feedback on the sparse SWMMAC fixes: - generate the adversarial A tile in-source (all six two-survivor position pairs, four single-survivor positions, zero-group; values keyed to (row, group, position)) and make it the DEFAULT build -- fails on the unfixed tree (max_abs_err 112/127/272 at K=32/64/128), passes after the fix; removes the uncommitted real_a_tile.h dependency - demote the slots-0,2 canonical pattern to -DUSE_CANONICAL_PATTERN control (it cannot detect the default bug) - scaffold a pk_int4 case behind -DENABLE_PK4_CASE, explicitly marked unvalidated (host fill layout unproven) pending an independent iu4 metadata-mapping cross-check - reword internal process labels to bug-number references; add ARCH SCOPE note at the packed path (the SWAP+XOR-1 mapping is gfx1201-measured) CCA

The-Monk · 2026-08-15T11:59:13Z

Thanks for this — running it on your own gfx1201 and bisecting the line both directions is exactly the review this needed, and the §5 finding (the canonical pattern structurally can't detect the bug) plus §1 (missing header) are both fair hits. Fixes pushed in ad79359:

§1/§5 — repro now self-contained, and the default config fails-before/passes-after. real_a_tile.h is gone; the test now generates its adversarial tile in-source: deterministic 2:4 input cycling every group through all six two-survivor position pairs, all four single-survivor positions, and the zero-survivor group, with values derived from (row, group, position) so misplacement shows numerically — same design as your substitute tile. That's the default build now. The old slots-0,2 synthetic pattern is demoted to an opt-in control (-DUSE_CANONICAL_PATTERN) with a comment stating it cannot detect the bug. Re-verified before pushing, same procedure you used: identical test source on both trees — head ALL PASS (max_abs_err 0 at K=32/64/128), base 8fc1ac2 FAIL (max_abs_err 112/127/272). So the buildable-by-default configuration is the one that fails on the unfixed tree.

§4 — iu4 coverage. Agreed this is the real gap, and I'll take you up on the one-hot v_swmmac_*_iu4 metadata sweep — that's the independent second measurement bug 3 needs, since my SWAP+XOR-1 mapping is device-measured but through my own transform stack. Say where you drop the generator/logs/disassembly and I'll cross-check against my derivation. Related question: do you have a working HIP-level iu4 route already — a __builtin_amdgcn_swmmac_*_iu4 builtin that lowers correctly on this clang, or an inline-asm wrapper from your probe work? Everything I have reaches the instruction through the CK pipeline; a bare intrinsic path would make the one-hot sweep trivially shareable and would also give my pk4 test an oracle route that doesn't touch the transform under test. On my side I've added a pk_int4 case scaffold to the test with an honest marker: the CPU oracle is straightforward (dense int4 reference from the same uncompressed logical values — not circular), but the host-side register-map fill convention for packed tensors is unproven on hardware, so it's compiled out by default (-DENABLE_PK4_CASE) until validated rather than pretending coverage.

§6 — intended scope. Bug 1's fix is intentionally architecture-generic: the old {a[2], a[3]} default is out-of-spec for any input whose survivors don't sit at 0/2, on every target the sparse selector reaches — your gfx942 compile-probe confirms the reach, and I'd argue the CDNA behavior change is the fix working as intended. If anyone with CDNA hardware can run the (now self-contained) repro there, great, but I don't think it should block. Bugs 2/3 are gfx12-only in practice today (no gfx9 pk_int4 sparse op, as you found); I've added an explicit ARCH SCOPE comment at the packed path stating the SWAP+XOR-1 mapping is gfx1201-measured and must be re-measured before any future CDNA pk4 sparse op trusts it.

§7 — housekeeping. Internal stage labels reworded to bug-number references throughout. Noted on the static_assert scope point — it's declaring the current implementation boundary, not narrowing behavior (nothing else reaches the packed branch today), and the scope comment now says so. Thanks for the add_subdirectory catch for the eventual gtest port — on the list for that change.

Your env notes (--rocm-device-lib-path, the libamdhip64.so symlink) match what I hit as well — I'll fold them into the test's header comment with the gtest port so the next person doesn't rediscover them.

The-Monk · 2026-08-15T12:07:35Z

Since you have gfx1201 and clearly care about this corner of the stack — some context on where this PR came from, and the wider map it belongs to. While building a 2:4-sparse fp8/int4 inference path for RDNA4 we kept hitting the same pattern: the capability is in the ISA, but no library exercises it, so the first real user finds the bugs (or the absence). The list of gaps we personally hit and had to fill with hand-written kernels, in case any of it is useful to you or worth more upstream issues:

Matrix/dot datapaths:

2:4-sparse SWMMAC (v_swmmac_*, ISA §7.12.3) — this PR's territory. CK's SPARSE family was silently wrong for arbitrary-position 2:4 (bug 1), the sparse selector is un-arch-specialized as you found, and as far as we can tell no inference stack used sparse prefill on RDNA4 at all before we did (our llama.cpp fork uses it for prefill: The-Monk/llama.cpp, roc8/roc10 branches).
iu4 packed sparse (v_swmmac_i32_16x16x{32,64}_iu4) — CK's packed path was structurally unreachable/broken (bugs 2/3), no gfx9 pk_int4 sparse op exists, the ISA's generic sparse pseudocode doesn't capture the iu4 idx semantics (hence the empirical SWAP+XOR-1), and we still don't know of a HIP builtin that lowers it correctly (the question upthread).
dp4a family for sub-2-bit quant (v_dot4_i32_iu8) — the generic HIP lowerings for 1-bit/ternary decode were poor on AMD (measured +37% on Q2_0 from one HIP-guarded vec_dot, cuda: AMD RDNA4 Q1_0/Q2_0 — HIP-path vec_dots (+37%/+16% decode), opt-in quant dedup, opt-in hipBLASLt prefill routes PrismML-Eng/llama.cpp#116) and binary had no dp4a decode at all; we wrote bit-spread+dp4a kernels for both.
FP8 (E4M3/E5M2) WMMA — mainline llama.cpp had no fp8 tensor types whatsoever; our fork added the types + kernels (bit-exact verified, ~96% of bandwidth roofline at 14B). MXFP8 similarly hardware-supported, zero stack coverage.

Verification gaps that compound it — validating new datapath kernels on gfx1201 is disproportionately hard because: device ASAN is unavailable (no xnack+ target-ID — exactly the tool that would have caught bug 1/2's OOB class), stochastic PC sampling is unsupported (host_trap only), and the GL2C EA size-split counters read zero even with full ppfeaturemask + profile_standard (base GL2C counters work — so it's counter-plumbing, not the perfmon clock). If you've gotten any of those three working on your 9070 XT I'd genuinely like to know how.

Happy to compare notes on any of these, or split out issues where they belong in other repos — the sparse/iu4 items are the CK-relevant ones, which is why they became this PR.

doplxyz · 2026-08-15T13:06:18Z

Re-ran everything on ad79359d0 on the 9070 XT. Short version: the new default
test reproduces exactly as you describe, and yes — there is a bare builtin
route to iu4, it runs on hardware, and it links no CK code.

All artifacts, logs, compile commands and disassembly for both rounds:
https://github.com/doplxyz/ck3759-gfx1201-verification @ 8cee1e3.

1. The bare iu4 intrinsic route you asked about — it exists

The builtins are on this toolchain. The name needs a _w32 suffix; without
it clang reports an undeclared identifier, which may be what you hit:

__builtin_amdgcn_swmmac_i32_16x16x32_iu4_w32(bool, int,  bool, int2, int8, int, bool);
__builtin_amdgcn_swmmac_i32_16x16x64_iu4_w32(bool, int2, bool, int4, int8, int, bool);

Seven arguments. I left the parameters unnamed above because I measured what
they do rather than assuming what they mean — one kernel per boolean
combination, read off the emitted asm:

(arg1, arg3, arg7)	modifiers on `v_swmmac_i32_16x16x32_iu4`
F, F, F	(none)
T, F, F	`neg_lo:[1,0,0]`
F, T, F	`neg_lo:[0,1,0]`
T, T, F	`neg_lo:[1,1,0]`
F, F, T	`clamp`
T, T, T	`neg_lo:[1,1,0] clamp`

So arg1 sets neg_lo[0], arg3 sets neg_lo[1], arg7 sets clamp, and
neg_lo[2] is not reachable from these arguments. CK passes arg1/arg3 as the A
and B signedness flags; I'm reporting the modifier mapping I measured rather
than re-asserting that interpretation. Probe source is
src/iu4_builtin_modifier_probe.cpp.

src/iu4_bare.cpp is a complete HIP program that includes nothing from
composable_kernel — no headers, no transforms, no layout helpers. It compiles,
launches, and returns data. Its device object contains exactly one iu4
instruction per kernel, zero iu8 instructions, and zero ck_tile
symbols:

v_swmmac_i32_16x16x32_iu4 v[0:7], v11, v[8:9], v12       neg_lo:[1,1,0]
v_swmmac_i32_16x16x64_iu4 v[0:7], v[12:13], v[8:11], v15 neg_lo:[1,1,0]

Per-lane operand sizes for wave32: K=32 takes A = 1 dword, B = 2 dwords,
D = 8 dwords; K=64 takes A = 2 dwords, B = 4 dwords, D = 8 dwords; idx is one
dword per lane in both. With every A and B nibble set to 1, every accumulator
dword in every lane reads 16 at K=32 and 32 at K=64 — the sparse product
counts — and a zero-A control returns all zeros. Bit-exact across fresh
processes.

Toolchain: AMD clang 23.0.0git, ROCm/llvm-project @ 46fcb339fb61, inside
rocm/pytorch:rocm7.14_ubuntu24.04_py3.12_pytorch_release_2.12.0. I have not
checked which older ROCm releases carry these builtins, so treat the _w32
names as confirmed for this snapshot only.

What this is and is not. It's a numeric route to the instruction that
doesn't pass through the transform under test. That's what makes a non-circular
second measurement possible, and it may be useful as an oracle route for your
pk4 case. It is not a measurement of idx semantics, and it is not a solved
register map: an all-ones stimulus is invariant under exactly the permutations
that matter, so it says nothing about which logical element landed where. I've
deliberately not formed a view yet on what any idx bit means. I'll pin your
current formula and its predicted table with a hash before I measure, so the
comparison can't drift after the fact.

To be precise about one thing, since you may want to reuse it: iu4_bare.cpp
doesn't avoid packed representation — the nibbles are still packed into dwords.
What it avoids is depending on CK's host-side fill helpers to do that packing,
because it writes and reads raw dwords directly. Establishing which raw nibble
position corresponds to which logical element is exactly the job of the
calibration below, not something the harness assumes.

2. `ad79359d0` re-verified — your numbers reproduce exactly

I used one immutable copy of your new test for every tree, so base and head
are never compared with different test sources. Clean builds (each variant's
output directory removed first), three fresh processes each. max_abs_err:

build	tree	K=32	K=64	K=128
default (adversarial)	base `8fc1ac2`	112	127	272	FAIL
default	head `ad79359d0`	0	0	0	PASS
`-DUSE_CANONICAL_PATTERN`	base `8fc1ac2`	0	0	0	PASS
`-DUSE_CANONICAL_PATTERN`	head	0	0	0	PASS
default	head minus bug 1	112	127	272	FAIL
default	base plus bug 1 only	0	0	0	PASS

All six are run1 == run2 == run3 byte-for-byte. Your 112/127/272 match to the
digit. §1 and §5 are closed from my side: the test is self-contained, the
default configuration is the one that fails on the unfixed tree, and the
canonical control passes on base — preserving the evidence that the old pattern
could not detect this.

Two further things from the same table. The head minus bug 1 build produces
stdout byte-identical to base, and bug 1 only passes — so for this test's
inputs, bug 1's one-line change is both necessary and sufficient for the
observed difference. That's a statement about this test, not about the fix's
full effect. And I disassembled all six builds: each contains 7
v_swmmac_i32_16x16x32_iu8 and zero iu4 instructions, so this table is a
bug-1 regression result and I'm not counting any of it as evidence for bugs 2
or 3.

I also checked that your header delta from ac24ac28d is comment-only: with
comments stripped, both sparse_transforms.hpp (163 lines) and
sparse_mma_pipeline.hpp (206 lines) are byte-identical across the two commits.
So my round-1 four-way bisect carries over to the current head without
re-running it. Hashes are in EVIDENCE_round2.md.

3. Nothing in the tree references the test

test/ck_tile/gfx1201_sparse_swmmac/ contains the .cpp and nothing else — no
CMakeLists.txt — and test/ck_tile/CMakeLists.txt has no add_subdirectory
for it. I grepped the whole head tree: the string gfx1201_sparse_swmmac
appears nowhere outside that directory. No CMake file, no script, no
Jenkinsfile stanza refers to it. So as the PR stands, nothing in the project's
own tooling can find this test — every result in this thread comes from
hand-driven builds, including mine.

I raised add_subdirectory last time as a footnote for the eventual gtest port.
That was too soft: whatever form the test finally takes, something in the tree
has to point at it, or the fix has no standing guard.

I want to be careful about the CI half of this, because I was about to overstate
it. The repository has no .github/ directory at all, so it isn't using GitHub
Actions and the empty check-suites on all three commits are the expected
consequence of that, not evidence of anything. CI here runs from the top-level
Jenkinsfile, which doesn't report into GitHub Checks — meaning I can't see from
outside whether it ran on this PR, what it covered, or whether it has gfx1201
hardware. That's worth asking a maintainer directly, and it's a different
question from the registration gap, which is visible in the tree and cheap to
close.

4. A thought on splitting bug 1 out

Entirely the maintainers' call and yours — I'm a drive-by reviewer with one GPU.

Bug 1 is in good shape: the test's default configuration fails before and passes
after on real hardware, independently reproduced, and the change is a plain
out-of-spec default. I find your §6 argument for keeping it
architecture-generic convincing. My probe result there is narrower than the
claim, though, and I should have said so last round: it shows that a gfx942
compile of the sparse path succeeds and emits v_smfmac_i32_16x16x64_i8. That
establishes the selector reaches gfx942 at compile time. It is not a runtime
result and doesn't generalize across gfx9 by itself.

Bugs 2 and 3 sit differently: the pk4 case is compiled out by default behind
-DENABLE_PK4_CASE, and the executing kernels are iu8-only, so as far as
anything in this thread shows, no number yet distinguishes bug 2 or bug 3 being
present from being absent.

Merging bug 1 now and keeping bugs 2 and 3 in a follow-up until the sweep and a
CK-level pk4 test land would get the confirmed fix into the tree without waiting
on the part that still needs measurement. Splitting inside this PR works just as
well if you'd rather not open another one — the point is only that the
verified part shouldn't have to wait on the unverified part.

5. What's next, and what I can't do

In order: the one-hot metadata sweep proper — full raw idx enumeration against
individually one-hot compressed slots, low and high nibble stimulated
separately, B basis-swept so I can tell which physical B element each result came
from, all 32 lanes and all 8 accumulator dwords collected, K=32 and K=64 swept
independently rather than one extrapolated from the other, with the lane/element
map established by calibration rather than borrowed from CK. Then the same
question end to end through CK's pk4 path, with bug-2-only and bug-3-only revert
mutants, since if those two can mask each other a head-vs-base comparison won't
separate them. Everything lands in the repo above with the raw tables.

Two things I can't close from here. I have no CDNA hardware, so my gfx942 result
stays compile-level; someone with a gfx942 part running the now-self-contained
repro would turn that into a runtime answer. And I can't tell whether this
repository's maintainers want a hardware-specific test registered in the normal
test graph, gated behind an option, or kept standalone — that decision, and what
evidence they'll accept for a target they may not have in CI, has to come from
them. Worth asking explicitly rather than either of us guessing.

On your verification-tooling questions (device ASAN, stochastic PC sampling, the
GL2C EA size-split counters reading zero): I have the same card and would rather
give you measurements than guesses, so I'll try all three and report back —
after the sweep, purely because I don't want to leave that half-finished. I've
subscribed to ROCm/ROCm#6613. I can see why you grouped them with this work; the
absence of device ASAN is precisely why bug 2's class had to be found the hard
way.

doplxyz · 2026-08-15T14:16:03Z

Sweep done, plus an end-to-end pk4 test. Headline: your bug-3 mapping is
correct, and I can now say why it has the shape it does. The run also turned up
something worth handling before this merges: the PR makes an existing test in
this repository fail.

Everything I measured is reproducible from
https://github.com/doplxyz/ck3759-gfx1201-verification @ b42b707. (Claims
below about what does or doesn't exist in the tree are from grepping it; claims
about history or anyone else's environment aren't things I can check.)

1. The one-hot iu4 metadata sweep — SWAP + XOR-1 confirmed

Method, as offered: bare __builtin_amdgcn_swmmac_*_iu4_w32, no
composable_kernel in the binary at all. For each (group, field pair, compressed slot) I drove one compressed A nibble and scanned every raw B position; a
nonzero accumulator names the position the hardware paired with that slot. No
logical coordinate system is assumed in the measurement — the grouping of B
positions and their order within a group come out of the data.

I pre-registered your rule and its predicted table, hashed and committed, before
taking any data (PREREGISTERED_HYPOTHESIS.md).

Result: your transform reproduces the hardware pairing exactly. 384 cells
across K=32 and K=64 — the full 4×4 field grid, both slots, every group, all 32
lanes — zero mismatches.

The raw law underneath:

idx field i governs compressed nibble i; a field carrying raw value v
pairs its nibble with the B nibble at raw offset v within that group.

Plain identity. No swap, no XOR. Both halves of your rule are the two places
CK_TILE_USE_PK4_LAYOUT_SHUFFLE enters — element 0 of a packed byte being the
HIGH nibble:

	CK index	raw index
compressed slot	`s` (0 = high)	`r = 1 - s` → the SWAP
uncompressed position	`j` (0 = byte0 high)	`o = j XOR 1` → the XOR

Substituting both into the identity law reproduces your code line for line, and
your corollary falls out: raw value 2 selects raw offset 2, which is CK position
3 — constant regardless of the real survivor, exactly the "always reconstructs at
position 3" you reported.

Adversarial checks, each breaking one way the first pass could have looked like
an identity without being one:

lane uniformity — one active A lane, every other lane carrying the opposite
idx: 1024/1024 followed the active lane's own idx. That rules out the specific
failure of another lane's metadata being substituted; it isn't a general proof
that no cross-lane routing exists.
arbitrary idx words — all groups live, a different random word per lane:
192/192.
superposition — leave-one-out on dense random inputs: 1919 nibbles, 0
mismatches.
operand range — the full signed 4-bit range including −8, with a non-zero
starting accumulator: 256/256 exact.
identical at -O0, -O2, -O3.
dataflow — between the operand loads and the instruction there are no
instructions at all touching the loaded values; the only ALU work in that
window is address arithmetic on the thread id. So the mapping isn't something
the lowering introduced.

A correction against my own case: my pre-registered transcription of your
rule was wrong. I wrote it in raw coordinates when it is defined in CK's,
silently assuming CK slot and position indices equal raw nibble indices. Scored
literally it "matches" 32/128, and anyone comparing the raw table against my
pre-registration would wrongly conclude your fix is broken. The 128/128 figure is
a post-hoc coordinate translation, not a confirmed prior prediction. Both are
recorded separately and the pre-registration is unmodified in git history.

On the ARCH SCOPE comment

Keep the instruction to re-measure — I was going to argue against it and I was
wrong. A future CDNA sparse op could differ in field assignment, operand
numbering or lane routing even with identical CK packing, so per-architecture
re-measurement stays necessary.

What I'd add is that the constants aren't encoding a gfx1201 oddity — the
gfx1201 encoding is the identity — they're encoding CK's packing convention.
Today that convention isn't switchable: config.hpp:181 defines
CK_TILE_USE_PK4_LAYOUT_SHUFFLE whenever it isn't already defined and
pk_int4.hpp tests it with #ifdef, so setting it to 0 changes nothing and
the #else branches there are unreachable. But that also means those constants
depend on a convention that a source change could flip without touching
architecture at all, and the #else branches are already dead code that would
silently disagree with the sparse path if revived. Worth a sentence in the same
comment.

2. `SparseTransformsTest.SingleNonZeroPerGroup` fails on this branch

This is the one I'd act on first.

test/ck_tile/core/arch/mma/pipeline/test_amdgcn_sparse_mma.cpp is registered
via _add_mma_gtest in test/ck_tile/core/arch/mma/CMakeLists.txt:17. It exists
at base 8fc1ac2 and this PR does not modify it. One of its cases drives
compress_a_impl on device with a single survivor per group and asserts:

// Single non-zero per group of 4 (at slot 3).
// nonzero_elems initializes to {a_vec[slot2]=0, a_vec[slot3]=V}.
// Only j=3 triggers: nonzero_elems[0]=V, field0=0b11, pos becomes 1.
// nonzero_elems[1] keeps its init V. Output: {V, V}.
expected_output[g * 2]     = val;
expected_output[g * 2 + 1] = val;

That expected value is bug 1's behaviour — the {a[2], a[3]} default reaching
the second compressed slot — recorded as the expected result.

I built and ran that file, unmodified, against both include trees, varying
nothing but -I:

include tree	result
base `8fc1ac2`	0 failing of 18
head `ad79359d0`	1 failing of 18 — `SparseTransformsTest.SingleNonZeroPerGroup`

EXPECT_EQ failed at test_amdgcn_sparse_mma.cpp:202   (x9, the compressed values)
compressed out, base: 5 5 6 6      <- what the test expects
compressed out, head: 5 0 6 0      <- what your fix produces
idx unchanged (0x000000bb) in both

MixedSparsityPattern, NonZerosAtSlots{1And3,0And3} and all eight
FullMatrixVerify_* cases pass on both trees; this is the only one that moves.

Caveat on method: googletest isn't installed in my container, so I supplied a
minimal gtest.h shim (TEST, EXPECT_*, ASSERT_*, GTEST_SKIP) rather than
building the CMake target. The test source itself is byte-for-byte upstream, and
the shim clearly does detect failures since it found this one — but I haven't run
the registered target through CMake, so treat this as "the upstream test's own
assertions fail on head", not "the CI job goes red" (which I can't observe
anyway).

Two things follow. The test needs updating as part of this PR, and that diff is
some of the better evidence in the change, since it shows the old behaviour
written down and corrected. And it explains why the two-survivor cases don't
move: bug 1 only bites when a group has fewer than two survivors, which
SingleNonZeroPerGroup is the only case in that file to supply. That's also the
shape of the gap new cases would fill.

It makes the registration point from my last comment concrete too: there's
already a registered home with the right idiom, so the standalone repro could
become extra cases in that file rather than a new directory — no new CMake entry
needed at all if it goes there.

3. End-to-end pk_int4 through CK

ENABLE_PK4_CASE expands to a printf and a TODO; it doesn't instantiate the
packed pipeline and asserts nothing. Grepping test/ and example/, the only
two files referencing the sparse pipeline are your standalone repro and
test_amdgcn_sparse_mma.cpp, and neither instantiates it with pk_int4_t. So as
far as the tree shows, bugs 2 and 3 have no numerical coverage in it, and the
numbers below are the first I can find evidence of for the packed path.

The fill convention you flagged as unproven, resolved. Asked CK's own types
rather than guessed:

for pk_int4_t the register map's vector index enumerates logical 4-bit
elements, not physical bytes — at K=32, num_vector_items is 16 while
sizeof(AWarpTensor) is 8, so porting the int8 fill verbatim overruns the
tensor by 2×;
logical element v goes to byte v/2, high nibble for even v;
the coordinate order is pinned without copying it from the code under test:
K ≠ N, so the two components have different ranges. Measured, A and B both
return coord[0] in [0,15] and coord[1] in [0,K-1], so B is
B[coord[1] * N + coord[0]]. I had this backwards at first and it produced a
plausible-looking wrong answer, which is why I went back and pinned it from the
ranges.

Logical k landing at raw nibble offset k XOR 1 is the same XOR the sweep
found from the hardware side. Not fully independent evidence — both involve the
same nibble convention — but they're arrived at from opposite ends.

Kill matrix. One independent build and run per shape, so a compile failure at
one shape can't mask another. max_abs_err against a dense int4 CPU reference
over the same logical values; int64 accumulation, exact comparison.

Every mutant row sits on top of the §4 fix, since without it nothing builds at
K≥128 and the rows would not be comparable. The first two rows show that fix's
effect on its own.

tree	K=32	K=64	K=128	K=256
head `ad79359d0`, unmodified	0 PASS	0 PASS	compile fail	compile fail
head + §4 fix — the baseline for the rows below	0 PASS	0 PASS	0 PASS	0 PASS
… + bug 2 reverted	0 PASS	0 PASS	compile fail	compile fail
… + bug 3 fully reverted	311 FAIL	451 FAIL	635 FAIL	904 FAIL
… + bug 3, swap kept, XOR removed	460 FAIL	810 FAIL	1142 FAIL	2290 FAIL
… + bug 3, XOR kept, swap removed	351 FAIL	588 FAIL	887 FAIL	1397 FAIL
… + bug 1's true-zero default reverted, packed path	64 FAIL	68 FAIL	89 FAIL	132 FAIL

Bug 3 is load-bearing, and in these configurations neither half alone
suffices — the CK-side counterpart to §1. (It doesn't rule out some other
equivalent formulation; only that these two partial forms are wrong.)
Bug 1's true-zero default is load-bearing in the packed path too, not only the
scalar path — the same defect §2's existing test records.
Bug 2's mutant survives at K=32 and K=64 and is killed only at K=128 and
K=256, as a compile error. The condition isn't "multi-fragment" as such but
"shapes where the idx word count diverges": at K=32/64 both accountings round
to a single word, so the difference is unobservable there. Whatever form a
regression test for bug 2 takes — end-to-end or a type-level unit test — it
needs a case where those counts differ, or it won't exercise that fix.

The test carries guard bands around every per-lane buffer, verified intact after
the fill so an off-by-a-factor is loud rather than silent; asserts stimulus
coverage rather than assuming it (all eleven at-most-two-survivor group patterns
present — the six two-survivor pairs, four single-survivor, one all-zero — and
the signed end point −8 present in both operands); and uses a B spanning the full
signed range, asymmetric in k and n so a transposed reading can't cancel out.

4. Multi-fragment packed A doesn't compile

At K=128 and K=256 the pk4 pipeline fails to build on head:

sparse_mma_pipeline.hpp:311: 'ATransformResult must match the return type of
ATransform::exec'   SparseIdxPack<2> vs SparseIdxPack<1>

checkATransformResult() re-derives the expected type as

decltype(ATransform::execExtVec(std::declval<ExternalAvecRef>()))

leaving LogicalADataType at its default. ext_vector_t<pk_int4_t, N> has
signed char as its scalar type, so the default resolves PackedSize to 1 —
while the real call path, exec(), passes ADataType explicitly and gets 2. The
check disagrees with the call it validates. It's hidden at K=32/64 because 8 and
16 two-bit fields both round to one idx word; the counts first diverge at K=128.

         static_assert(
             std::is_same_v<ATransformResult,
-                           decltype(ATransform::execExtVec(std::declval<ExternalAvecRef>()))>,
+                           decltype(ATransform::template execExtVec<AVecType, ADataType>(
+                               std::declval<ExternalAvecRef>()))>,
             "ATransformResult must match the return type of ATransform::exec");

With that, all four shapes compile and pass, and your int8 tests are unaffected
(adversarial ALL PASS, canonical control ALL PASS). Whether K≥128 is a shape the
packed path is meant to support is your and the maintainers' call — the argument
that it is comes only from your int8 test covering K=128 and the pipeline being a
template over K, which isn't conclusive for a different data type. What is
measured is that the check and the call disagree. I haven't run a compile matrix
over every type/shape combination, so read the diff as "makes the check agree
with the call", not as "verified side-effect-free".

5. What this doesn't show

Mutant kills show the test is sensitive to the code under test; they aren't proof
head is correct. If head and my harness shared a wrong layout assumption, mutants
could still die while a correlated error passed — the range-based determination
of the coordinate order and the CK-free sweep are what reduce that risk, not the
kill matrix. Everything is FragsM = FragsN = 1, gfx1201, one toolchain, wave32,
full EXEC. I still have no CDNA hardware, so bug 1's reach there stays a
compile-level result on my side.

Happy to open a PR against your branch with any of this rather than leave you to
lift it — the pk4 case and the SingleNonZeroPerGroup update being the two that
matter most for someone reviewing without a gfx1201 part.

…e2e; fix pk4 K>=128 type check Review-round changes, all four from doplxyz's independent verification (github.com/doplxyz/ck3759-gfx1201-verification): - SparseTransformsTest.SingleNonZeroPerGroup expected {V, V}: the second compressed slot leaking a_vec[slot3] through the pre-fix {a[2], a[3]} default -- bug 1 recorded as the expected result. Corrected to {V, 0} with a HISTORY note; this was the only registered case supplying a group with fewer than two survivors, which is exactly where bug 1 bites. - The standalone repro moves from test/ck_tile/gfx1201_sparse_swmmac/ (unregistered, no CMakeLists) into test_amdgcn_sparse_mma.cpp as SparseSwmmacE2E.{AdversarialGeneratedTile,CanonicalPatternControl}; the compile-time USE_CANONICAL_PATTERN toggle becomes a runtime parameter so the control runs in the same binary. Runtime-skipped on non-gfx12 devices. - New SparsePk4E2E.AdversarialInt4AllShapes (K=32/64/128/256): first in-tree numerical coverage for the packed-nibble path (bugs 2/3). Ported with attribution from doplxyz's MIT-licensed pk4_e2e harness: dense int4 CPU oracle over logical values (independent of the transform under test), guard-banded host fill via CK's own register maps. K=128/256 also lock bug 2 and the type-check fix below at compile time -- idx word counts only diverge at FragsK > 1. - checkATransformResult (sparse_mma_pipeline.hpp) re-derived the expected transform type via execExtVec with the defaulted LogicalADataType, resolving PackedSize to 1 while the real exec call passes ADataType and gets 2 -- so the pk4 pipeline failed to compile at K>=128. Fix (explicit template arguments) by doplxyz, applied verbatim. - ARCH SCOPE comment extended: one-hot sweeps with bare v_swmmac_*_iu4_w32 builtins show the gfx1201 idx law is a plain identity in raw nibble coordinates; SWAP + XOR-1 encode CK_TILE_USE_PK4_LAYOUT_SHUFFLE's high-nibble-first convention, not a hardware quirk, and config.hpp pins that macro on (the pk_int4.hpp #else branches are unreachable). Convention dependency documented alongside the existing per-architecture re-measurement instruction. Verified on gfx1201 (Radeon AI PRO R9700, ROCm 7.14): all 21 tests in test_amdgcn_sparse_mma pass, including the four pk4 shapes (max_abs_err = 0 exact), the adversarial e2e that fails on the unfixed tree, and the corrected single-survivor case. CCA

The-Monk · 2026-08-16T08:19:26Z

This is a remarkable verification round — the pre-registration discipline, the adversarial follow-ups on your own first-pass results, and the coordinate-system self-correction you flagged against your own case are all above and beyond. Everything actionable from your three comments is now in the branch as 8c60a1249. Point by point:

1. SingleNonZeroPerGroup — fixed, and you're right that the diff is the evidence. The expectation is corrected to {V, 0} with a HISTORY comment stating plainly that the old {V, V} recorded bug 1's leaked-survivor dataflow as correct behavior. Your observation about why the two-survivor cases don't move (bug 1 only bites below two survivors) is captured there too.

2. Repro folded into the registered file. test/ck_tile/gfx1201_sparse_swmmac/ is gone; the standalone repro now lives in test_amdgcn_sparse_mma.cpp as SparseSwmmacE2E.AdversarialGeneratedTile (K=32/64/128) plus SparseSwmmacE2E.CanonicalPatternControl — the control kept, as you put it, as preserved evidence that the legacy pattern cannot detect the default bug. The compile-time USE_CANONICAL_PATTERN toggle became a runtime parameter so both run in one binary. No new CMake entry, exactly as you suggested.

3. pk4 numerical coverage — ported from your harness, with attribution. Your pk4_e2e.cpp resolved the two things that kept me from writing this test honestly (the logical-nibble register map convention, and the (n, k) coordinate order pinned from ranges rather than read off the code under test), and it's MIT-licensed, so rather than re-derive a worse version I ported it into the file as SparsePk4E2E.AdversarialInt4AllShapes (K=32/64/128/256) with an attribution comment pointing at your repo. That gives bugs 2 and 3 their first in-tree numerical coverage, and per your point about idx-word divergence, the K=128/256 cases also lock bug 2 and the §4 fix at compile time — a revert of either fails the build before it can fail the assert. If you'd rather the port be structured differently (or prefer to contribute it under your own name in a follow-up), happy to restructure; the attribution stands either way.

4. §4 (checkATransformResult) — your patch applied verbatim, credited in the commit. The diagnosis was exact: the check re-derived the transform type with the defaulted LogicalADataType (PackedSize 1) while the real call passes ADataType (PackedSize 2). With it, head compiles and passes at K=128/256 where it previously couldn't build at all.

5. ARCH SCOPE comment extended with your convention-scope point. It now states that the measured hardware idx law is a plain identity in raw nibble coordinates, that SWAP + XOR-1 are the coordinate change from CK_TILE_USE_PK4_LAYOUT_SHUFFLE's high-nibble-first convention (not a gfx1201 quirk), and that config.hpp defines the macro unconditionally so pk_int4.hpp's #else branches are unreachable dead code that must not be revived without re-deriving these constants. The per-architecture re-measurement instruction stays, per your argument for keeping it.

On the _w32 builtins — that was exactly my miss. I had probed the _gfx12-suffixed spellings and concluded no builtin existed; your modifier-mapping table (arg1/arg3 → neg_lo[0]/[1], arg7 → clamp) and the bare-route program settle it, and the non-circular oracle route is now part of my toolbox. Related: while debugging WMMA fragment layouts this week I posted an independent lane-mapping verification (identity-matrix probe, 256/256) to ROCm/ROCm#6025 — the f16 case of the same documentation gap your one-hot method addresses for the sparse idx semantics.

Results on gfx1201 (Radeon AI PRO R9700, ROCm 7.14), full registered suite:

[----------] SparseSwmmacE2E
[K32_single_frag] M=16 N=16 K=32 -> max_abs_err=0 PASS
[K64_2frag]      M=16 N=16 K=64 -> max_abs_err=0 PASS
[K128_4frag]     M=16 N=16 K=128 -> max_abs_err=0 PASS
[K32_control]    M=16 N=16 K=32 -> max_abs_err=0 PASS
[K64_control]    M=16 N=16 K=64 -> max_abs_err=0 PASS
[----------] SparsePk4E2E
[pk4_K32]  M=16 N=16 K=32  -> max_abs_err=0 PASS
[pk4_K64]  M=16 N=16 K=64  -> max_abs_err=0 PASS
[pk4_K128] M=16 N=16 K=128 -> max_abs_err=0 PASS
[pk4_K256] M=16 N=16 K=256 -> max_abs_err=0 PASS

[==========] 21 tests from 5 test suites ran. (104 ms total)
[  PASSED  ] 21 tests.  (incl. the corrected SingleNonZeroPerGroup)

One honest caveat on portability: the new E2E cases GTEST_SKIP at runtime on non-gfx12 devices, and instantiate Gfx1201Target pipelines unconditionally at compile time — same pattern as this file's existing gfx950-target cases, relying on CK's internal per-arch guards for foreign device passes. If a multi-target CI build objects, the fix is a coarse #if around the two namespaces; I didn't pre-add it since the gfx950 precedent suggests it's unnecessary.

doplxyz · 2026-08-18T14:16:40Z

Re-verified 8c60a1249 on my own gfx1201. Every production fix site covered by the matrix is pinned by a
test failure or a compile failure, so no mutant survives. I also found a build regression that I think blocks the merge,
and a possible CMake/CI gate mismatch that I can only half-verify. Both come with measurements, and
the regression comes with a patch I have tested.

1. On the pk4 port

Please keep it as it is. I wrote that harness and published it under MIT; I am happy for it
to be included under CK's project licensing, and I do not need a separate authorship commit
or a restructuring. No third-party code is involved. Landing the coverage in this PR is
worth more to me than a separate commit under my name.

2. What I tested


Tested commit	`8c60a1249da1bb58262ac29e143ce59beb4dff6e` (PR head)
Base	`f7982554ee23f9f5e3d12d3c7bbeea942c988bf8`
GPU	Radeon RX 9070 XT (gfx1201)
ROCm	7.2.4 (`/opt/rocm/.info/version`; `hipconfig --version` → `7.2.53211-97f5574fe2`)
Compiler	AMD clang 22.0.0git `roc-7.2.4`
Build	CMake 4.4.2 + Ninja, real GoogleTest via the repo's `FetchContent`, target `test_amdgcn_sparse_mma`

The ROCm version differs from yours (7.14.0), so this is a cross-version reproduction rather
than a repeat of your run.

With GPU_TARGETS=gfx1201, clean configure and build:

[==========] 21 tests from 5 test suites ran. (447 ms total)
[  PASSED  ] 21 tests.

--gtest_list_tests lists all 21 including the three new ones, nothing was skipped, and HIP
reported the device as gfx1201 — so the new E2E cases really executed rather than silently
skipping.

3. Mutation matrix

Test code held at head, one production edit per mutant, the target's objects deleted and
rebuilt for every row. Both headers verified byte-identical to head afterwards.

mutant	kind	build	killed by
baseline	—	ok	(21/21 pass)
bug 1, scalar path (`nonzero_elems` default)	historical revert	ok	`SparseTransformsTest.SingleNonZeroPerGroup` + `SparseSwmmacE2E.AdversarialGeneratedTile`
bug 1, packed path (`survivor` default)	synthetic	ok	`SparsePk4E2E.AdversarialInt4AllShapes`
bug 2 (`* MmaOp::APackedSize`)	historical revert	fails	`static_assert`: `SparseIdxPack<2>` vs `<1>` (`sparse_mma_pipeline.hpp:326`)
bug 3, SWAP + XOR	synthetic	ok	`SparsePk4E2E.AdversarialInt4AllShapes`
bug 3, SWAP only	synthetic	ok	same
bug 3, XOR only	synthetic	ok	same
`checkATransformResult` type fix	historical revert	fails	`static_assert`: `ATransformResult` (`sparse_mma_pipeline.hpp:311`)

No mutant survived.

The corrected SingleNonZeroPerGroup expectation is load-bearing. Reverting only the
scalar-path default makes it fail. I want to correct my own earlier framing here: I had
described this as "passes on base and passes on head", which is not evidence — those are two
different oracles. The mutant is the evidence.

Bug 2 is guarded at compile time, not numerically. Even with the checkATransformResult
type fix in place, reverting bug 2 trips a different static_assert before anything runs,
so the numerical consequence of that bug is still unobserved. Your claim that K=128/256 lock
it at compile time holds, and the lock is independent of the type fix — but the suite gives
bug 2 no numerical coverage. That seems fine to me (a compile-time lock is harder to ignore
than a failing assert); it just should not be described as numerical coverage.

The packed-path rows are synthetic mutants — that path does not exist at base — so they show
test sensitivity, not a regression that ever shipped.

I also checked whether more shapes were worth adding: sparse_mma_pipeline.hpp static_asserts
FragsM == 1 and FragsN == 1, so those shapes are unreachable by construction. FragsK is
the only fragment dimension that varies, and K=32/64/128/256 already covers 1/2/4/8. I do not
think any additional full-EXEC fragment shape is missing.

4. The `test_amdgcn_sparse_mma` target no longer builds for gfx9 targets

This is the one I would treat as blocking. I only built this target, so this is a statement
about it and not about the whole tree.

`GPU_TARGETS`	base `f7982554`	head `8c60a1249`
gfx942	builds	fails
gfx950	builds	fails
gfx942;gfx1201	builds	fails
gfx12-generic;gfx1201	builds	builds
gfx1201	builds	builds

include/ck_tile/core/arch/mma/sparse/wmma/sparse_gfx12.hpp:153:17: error:
  '__builtin_amdgcn_swmmac_i32_16x16x32_iu8_w32' needs target feature gfx12-insts,wavefrontsize32
  ... also :287 (iu4 16x16x32) and :313 (iu4 16x16x64)

The gfx942 and gfx950 stages are both ON by default in Jenkinsfile.

I think the gfx950 precedent did not transfer because it is a different kind of test.
SparseMMATrait.SparseMfmaGfx950Specialization is a pure type-trait check — it names
CompilerTargetGfx950 but never instantiates a pipeline or emits a builtin. The existing
pipeline tests (FullMatrixVerify_*) go through mma_pipeline_test::run_pipeline_matrix_test,
which dispatches over the configured CMake targets. The new E2E cases name Gfx1201Target
directly and bypass both, so the gfx12 SWMMAC wrappers are instantiated in every device pass,
including gfx9 ones. At base nothing instantiated them, which is why base builds.

I first tried guarding this in the test with a compile-time counterpart to device_is_gfx12()
driven by CK_CMAKE_GPU_TARGET_IDS. That fixes single-target gfx9 builds but not
gfx942;gfx1201, because that macro lists all configured targets and is identical in every
device pass. Guarding the wrappers themselves is what actually works, and it matches the
existing #if defined(__gfx950__) idiom in unary_element_wise_operation.hpp:

     exec(AVecType const& aVec, BVecType const& bVec, CVecType const& cVec, int32_t idx)
     {
+        #if defined(__GFX12__)
         using P = WarpGemmParamsParser<Params...>;
         return {__builtin_amdgcn_swmmac_i32_16x16x32_iu8_w32(true, // A signedness
                                                              aVec,
@@
                                                              P::clamp)};
+        #else
+        // Not a gfx12 device pass: this specialization can be instantiated in a
+        // multi-target build (e.g. GPU_TARGETS="gfx942;gfx1201"), where the
+        // builtin is unavailable. Unreachable as long as dispatch only calls
+        // this on a gfx12 device; a mis-dispatch traps rather than miscomputes.
+        (void)aVec;
+        (void)bVec;
+        (void)cVec;
+        (void)idx;
+        __builtin_trap();
+        __builtin_unreachable();
+        #endif
     }

applied to the three wrappers the new tests instantiate (iu8 16x16x32, iu4 16x16x32, iu4
16x16x64). Two details -Werror cares about: using P has to move inside the #if
(-Wunused-local-typedef), and the parameters need consuming in the #else
(-Wunused-parameter). The other eight wrappers in that header would need the same treatment if anything
instantiates them from a multi-target build; nothing in the configurations I built does, but
I have not checked the whole tree. If the intent is to make the header generally
mixed-target safe rather than just fix this regression, all eleven should be guarded
consistently.

With that patch:

`GPU_TARGETS`	build	run
gfx942	ok	—
gfx950	ok	—
gfx942;gfx1201	ok	21/21 pass
gfx12-generic;gfx1201	ok	21/21 pass
gfx1201	ok	21/21 pass

Note this keeps gfx12 coverage in mixed builds, which the CK_CMAKE_GPU_TARGET_IDS approach
I tried would have thrown away. I did not try a test-side guard driven by __GFX12__. I can open this as a PR against your branch if that is easier, or you can take the diff.

5. A CMake gate that may exclude this file from the gfx1201 CI job

Smaller, and I can only verify half of it.

The target is guarded by test/ck_tile/core/arch/mma/CMakeLists.txt:16:

if(GPU_TARGETS MATCHES "gfx9|gfx120")

Measured with ninja -t targets all:

`GPU_TARGETS`	configure	`test_amdgcn_sparse_mma`
gfx1201	ok	generated
gfx12-generic	ok	not generated
gfx942 / gfx950	ok	generated
gfx1250	ok	not generated

gfx12-generic matches neither alternative. What makes me think this matters is that
pipeline_tests_helper.hpp already anticipates exactly that configuration:

// gfx12-generic and gfx11-generic make no difference with the specialized archs.
// Some CI pipelines make use of that and configure the project with the generic
// flags besides compiling for (f.e.) gfx1201.

So the runtime side expects a gfx12-generic build running on a gfx1201 device, while the
CMake gate excludes that configuration outright. A job configured that way would configure
successfully and quietly build no sparse-MMA test target at all — a silent omission rather
than a failure. Jenkinsfile:843 passes "gfx12-generic" to ck.runBuildCKAndTests on a
rocmnode("gfx1201") agent, but that helper lives in the external ck shared library, so I
cannot see what reaches -DGPU_TARGETS.

Could someone with access check that job's CMakeCache.txt or configure log for the effective
GPU_TARGETS, and whether test_amdgcn_sparse_mma appears in its targets? If it does need
fixing, I would add the entry rather than broaden the regex to gfx12, since gfx12 as a
regex also matches gfx1250:

if(GPU_TARGETS MATCHES "gfx9|gfx120" OR "gfx12-generic" IN_LIST GPU_TARGETS)

I verified that this builds and lists all three new tests under GPU_TARGETS=gfx12-generic
and changes nothing for gfx1201, gfx942 or gfx1250.

One related observation: device_is_gfx12() uses strstr(gcnArchName, "gfx12"), which also
matches gfx1250, and amdgcn_target_id::GFX1250 is 0x1250, which falls inside the
GFX1200 .. GFX12_GENERIC (0x1200 .. 0x12FF) range that pipeline_tests_helper.hpp uses
for its generic-target handling. Neither matters today, but both would if gfx1250 ever built
this file.

6. Limits

gfx942 and gfx950 are compile-only here; I have no CDNA hardware, so nothing about
runtime behaviour on those parts is verified.
Partial EXEC is not tested. I could not find anything in sparse_mma_pipeline.hpp
stating that a full wave is a precondition. Since these fixes touch survivor selection and
nibble shuffling, could you confirm whether partial EXEC is simply unsupported for this
primitive? If so, a one-line comment saying that would close it.
gfx1250 is unevaluated: it does not build in my ROCm at all, for an unrelated reason
(hipLaunchAttributeClusterDimension is absent in 7.2.4).
None of this substitutes for a CI run.

I can provide the exact configure/build commands, the per-mutant diffs, and full logs for any
row above.

@doplxyz

…l-wave requirement The gfx12 SWMMAC builtin wrappers (iu8 16x16x32, iu4 16x16x32/64) are instantiated by the new E2E tests via Gfx1201Target directly, so in a multi-target build (e.g. GPU_TARGETS=gfx942;gfx1201, both ON by default in Jenkinsfile) the gfx12 builtins are emitted in gfx9 device passes where they do not exist, breaking the test_amdgcn_sparse_mma build. Guard the three instantiated wrappers with #if defined(__GFX12__), matching the existing __gfx950__ idiom; the non-gfx12 pass traps rather than miscomputes on any mis-dispatch. Also document that this primitive requires a full active wavefront (per-lane fragment distribution), so partial EXEC is unsupported. Reported and patch-tested by @doplxyz on RX 9070 XT / ROCm 7.2.4.

@doplxyz

test/.../mma/CMakeLists.txt gated test_amdgcn_sparse_mma on GPU_TARGETS MATCHES gfx9|gfx120, which gfx12-generic does not match, so a CI job configured with gfx12-generic on a gfx1201 device (as pipeline_tests_helper.hpp anticipates) would silently build no sparse-MMA test target. Add an explicit gfx12-generic term rather than broadening to gfx12 (which also matches gfx1250). Reported by @doplxyz.

The-Monk · 2026-08-18T15:50:23Z

@doplxyz — thank you, this is an exceptional review. The cross-version reproduction on 7.2.4 and the full mutation matrix are more than I could have asked for, and you're right that the mutant is the evidence, not the base-vs-head pass. I've pushed ddb8567f1 addressing the blocker; point-by-point below.

Build regression (gfx9 / multi-target) — fixed

Applied your #if defined(__GFX12__) guard verbatim to the three wrappers the E2E tests instantiate (iu8 16x16x32, iu4 16x16x32, iu4 16x16x64), including the two -Werror details (using P moved inside the #if, parameters consumed in the #else). The non-gfx12 pass traps rather than miscomputes on a mis-dispatch, matching the __gfx950__ idiom you pointed at. Credited you in the commit.

On the other eight wrappers: I scoped the guard to the three actually instantiated in a device pass today, to keep the change to what's tested. I agree the header should be mixed-target-safe as a property, not just for these three — if you'd prefer I extend the same guard to all eleven in this PR, say the word and I will (it's mechanical; I just didn't want to ship guards on wrappers neither of us has build-exercised in a mixed target).

Partial EXEC — unsupported, now documented

Confirmed: this primitive requires a full active wavefront. The A/B/C fragments are distributed per-lane across all 32 lanes (kABKPerLane in sparse_mma_pipeline.hpp), so the swmmac builtins are wave-cooperative — partial EXEC leaves those fragment lanes undefined. I added a one-line note to the pipeline doc block in ddb8567f1. That's why the survivor-selection and nibble-shuffle fixes never had to consider a partial wave.

Bug-2 framing — agreed

You're correct that bug 2 is guarded at compile time (the SparseIdxPack<2> vs <1> static_assert) and gets no numerical coverage, independent of the type fix. I won't describe it as numerical coverage. As you say, a compile-time lock is harder to ignore than a failing assert, so I'm comfortable leaving it there — but the distinction is fair and I appreciate you drawing it precisely.

CMake / CI gate (`gfx12-generic`) — I think you're right, and I can't see the CI either

I can confirm the local half: test/ck_tile/core/arch/mma/CMakeLists.txt:16 is if(GPU_TARGETS MATCHES "gfx9|gfx120"), and gfx12-generic matches neither alternative, so the target is silently not generated under that configuration — while pipeline_tests_helper.hpp explicitly anticipates a gfx12-generic build running on a gfx1201 device. I don't have visibility into what ck.runBuildCKAndTests passes as -DGPU_TARGETS on the rocmnode("gfx1201") agent, so I can't confirm whether the real CI job hits this.

I've folded your fix into this PR (84fbf0de5) rather than broaden to gfx12 (which would also match gfx1250):

if(GPU_TARGETS MATCHES "gfx9|gfx120" OR "gfx12-generic" IN_LIST GPU_TARGETS)

so the target builds under either configuration regardless of what the CI job turns out to pass; if someone with CI access confirms the effective GPU_TARGETS on the rocmnode("gfx1201") agent we can revisit, but this makes it robust either way. The related device_is_gfx12() / gfx1250-range observations are good catches for whenever gfx1250 builds this file; I'll leave a note rather than change behavior today.

Thanks again — this is the most thorough external validation this work has had, and the patch saved me the mixed-target debug.

doplxyz · 2026-08-18T16:34:24Z

Unrelated to this PR's review — just closing the loop on the offer I made earlier. I ran the three
RDNA4 verification-tooling items you listed on my 9070 XT. Short answer to "if you've gotten any of
those three working": none of them, but two came with a sharper description than "unsupported", and
two of my own first-pass conclusions turned out to be wrong and are corrected in the write-up.

Posted as a comment on your ROCm/ROCm#6613 rather than as new issues, since it is your ticket and
splitting it seemed worse than keeping it in one place:
ROCm/ROCm#6613 (comment)

Measurements, scripts and raw logs:
https://github.com/doplxyz/rdna4-gfx1201-tooling-verification

Headline for each: device ASAN — every RDNA target from gfx1030 on rejects xnack+, not just
RDNA4, and on gfx1201 -fgpu-sanitize emits a device .text byte-identical to the plain build
while exiting 0. Stochastic PC sampling — not advertised, rejected in all 12 configurations tried,
with host_trap working on the same workload as a control. GL2C — the base EA counters are exact
(count = bytes/256 + 2 across a size sweep, so 256 B per request), the size-split counters are
zero on every run that lands on that fit, and FETCH_SIZE is defined solely from those three, so
it reports 0 KB for a kernel that read 256 MiB. Same on ROCm 7.2.4 and 7.14.

The ball on this PR is still yours; nothing here needs a reply from you.

yraparti and others added 30 commits May 29, 2026 17:09

The-Monk requested review from Snektron and vpietila-amd as code owners July 24, 2026 03:03

The-Monk mentioned this pull request Jul 24, 2026

ck_tile 2:4-sparse SWMMAC (MmaOpFamily::SPARSE) produces incorrect results on gfx1201/RDNA4 — 3 bugs in compress_a_impl / iu4 path #3753

Open

arai713 and others added 6 commits July 24, 2026 17:29

Merge branch 'develop' into gfx1201-sparse-swmmac-fixes

ac24ac2

The-Monk mentioned this pull request Aug 5, 2026

Track 2, Justin Monk, Hyperloom AMD-DEV-CONTEST/Radeon-hackathon-2026-07#99

Open

The-Monk requested a review from JiaLuo-CAN as a code owner August 15, 2026 11:50

The-Monk mentioned this pull request Aug 15, 2026

[Feature]: Device-side ASAN (or equivalent OOB detection) for RDNA4 — gfx1201 rejects xnack+ target-ID ROCm/ROCm#6613

Open

The-Monk force-pushed the gfx1201-sparse-swmmac-fixes branch from 942eb11 to 8c60a12 Compare August 16, 2026 08:15

jmonk added 2 commits August 18, 2026 11:40

shumway force-pushed the develop branch from 8829582 to 24fd85f Compare August 18, 2026 18:23

shumway requested review from a team and ddembeckAMD as code owners August 18, 2026 18:23

Navigation Menu

Uh oh!

ck_tile: fix 2:4-sparse SWMMAC correctness on gfx1201/RDNA4 (3 bugs) + fail→pass repro - #3759

ck_tile: fix 2:4-sparse SWMMAC correctness on gfx1201/RDNA4 (3 bugs) + fail→pass repro#3759
The-Monk wants to merge 144 commits into
ROCm:developfrom
The-Monk:gfx1201-sparse-swmmac-fixes

The-Monk commented Jul 24, 2026 •

edited

Loading

Uh oh!

doplxyz commented Aug 14, 2026

Uh oh!

The-Monk commented Aug 15, 2026

Uh oh!

The-Monk commented Aug 15, 2026

Uh oh!

doplxyz commented Aug 15, 2026

Uh oh!

doplxyz commented Aug 15, 2026

Uh oh!

The-Monk commented Aug 16, 2026

Uh oh!

doplxyz commented Aug 18, 2026

Uh oh!

The-Monk commented Aug 18, 2026

Uh oh!

doplxyz commented Aug 18, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

20 participants

Uh oh!

Conversation

The-Monk commented Jul 24, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

0. TL;DR for the maintainer (schung-amd)

1. The three bugs (files: sparse_mma_pipeline.hpp, sparse_transforms.hpp)

2. The test (committed; fails on current CK, passes with the fix)

3. Context — why these bugs mattered (the "whole thing")

4. Closing the gap — result

5. A related RDNA4 iu4 quirk (same class, different instruction — FYI, not part of the fix)

6. Method (for reproducibility)

Uh oh!

doplxyz commented Aug 14, 2026

1. The repro can't be built from the PR alone

2. base fails, head passes

3. The whole effect is bug 1

4. The test never reaches iu4, so bugs 2 and 3 get no numerical coverage from it

5. The buildable configuration passes on the unfixed tree — and the reason is the one you identified

6. The change isn't scoped to gfx1201, and bug 1's fix demonstrably reaches CDNA

7. Minor, non-blocking

What would let me say more

Uh oh!

The-Monk commented Aug 15, 2026

Uh oh!

The-Monk commented Aug 15, 2026

Uh oh!

doplxyz commented Aug 15, 2026

1. The bare iu4 intrinsic route you asked about — it exists

2. ad79359d0 re-verified — your numbers reproduce exactly

3. Nothing in the tree references the test

4. A thought on splitting bug 1 out

5. What's next, and what I can't do

Uh oh!

doplxyz commented Aug 15, 2026

1. The one-hot iu4 metadata sweep — SWAP + XOR-1 confirmed

On the ARCH SCOPE comment

2. SparseTransformsTest.SingleNonZeroPerGroup fails on this branch

3. End-to-end pk_int4 through CK

4. Multi-fragment packed A doesn't compile

5. What this doesn't show

Uh oh!

The-Monk commented Aug 16, 2026

Uh oh!

doplxyz commented Aug 18, 2026

1. On the pk4 port

2. What I tested

3. Mutation matrix

4. The test_amdgcn_sparse_mma target no longer builds for gfx9 targets

5. A CMake gate that may exclude this file from the gfx1201 CI job

6. Limits

Uh oh!

The-Monk commented Aug 18, 2026

Build regression (gfx9 / multi-target) — fixed

Partial EXEC — unsupported, now documented

Bug-2 framing — agreed

CMake / CI gate (gfx12-generic) — I think you're right, and I can't see the CI either

Uh oh!

doplxyz commented Aug 18, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

20 participants

The-Monk commented Jul 24, 2026 •

edited

Loading

1. The three bugs (files: `sparse_mma_pipeline.hpp`, `sparse_transforms.hpp`)

2. `ad79359d0` re-verified — your numbers reproduce exactly

2. `SparseTransformsTest.SingleNonZeroPerGroup` fails on this branch

4. The `test_amdgcn_sparse_mma` target no longer builds for gfx9 targets

CMake / CI gate (`gfx12-generic`) — I think you're right, and I can't see the CI either